Remove fault_tolerance header and macros

johnbowen42 · johnbowen42 · commit 1f93adf93df3 · 2025-12-22T09:54:29.000-08:00
diff --git a/include/RAJA/internal/fault_tolerance.hpp b/include/RAJA/internal/fault_tolerance.hpp
diff --git a/include/RAJA/policy/cuda/WorkGroup/WorkRunner.hpp b/include/RAJA/policy/cuda/WorkGroup/WorkRunner.hpp
@@ -377,7 +377,6 @@ struct WorkRunner<
                                          block_size),
           static_cast<cuda_dim_member_t>(num_loops), 1};
 
-      RAJA_FT_BEGIN;
 
       //
       // Setup shared memory buffers
@@ -393,7 +392,6 @@ struct WorkRunner<
                            shmem, r, Async);
       }
 
-      RAJA_FT_END;
     }
 
     return run_storage;
diff --git a/include/RAJA/policy/cuda/forall.hpp b/include/RAJA/policy/cuda/forall.hpp
@@ -570,7 +570,6 @@ forall_impl(resources::Cuda cuda_res,
     internal::CudaDims dims(1);
     DimensionCalculator::set_dimensions(dims, len, func, shmem);
 
-    RAJA_FT_BEGIN;
 
     RAJA::cuda::detail::cudaInfo launch_info;
     launch_info.gridDim  = dims.blocks;
@@ -599,7 +598,6 @@ forall_impl(resources::Cuda cuda_res,
                                                       launch_info);
     }
 
-    RAJA_FT_END;
   }
 
   return resources::EventProxy<resources::Cuda>(cuda_res);
diff --git a/include/RAJA/policy/cuda/launch.hpp b/include/RAJA/policy/cuda/launch.hpp
@@ -96,7 +96,6 @@ struct LaunchExecute<
         blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       size_t shared_mem_size = launch_params.shared_mem_size;
       RAJA::cuda::detail::cudaInfo launch_info;
@@ -127,7 +126,6 @@ struct LaunchExecute<
                                                         launch_info);
       }
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
@@ -207,7 +205,6 @@ struct LaunchExecute<
         blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       size_t shared_mem_size = launch_params.shared_mem_size;
       RAJA::cuda::detail::cudaInfo launch_info;
@@ -238,7 +235,6 @@ struct LaunchExecute<
                                                         launch_info);
       }
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
diff --git a/include/RAJA/policy/hip/WorkGroup/WorkRunner.hpp b/include/RAJA/policy/hip/WorkGroup/WorkRunner.hpp
@@ -364,7 +364,6 @@ struct WorkRunner<
                                         block_size),
           static_cast<hip_dim_member_t>(num_loops), 1};
 
-      RAJA_FT_BEGIN;
 
       //
       // Setup shared memory buffers
@@ -380,7 +379,6 @@ struct WorkRunner<
                           shmem, r, Async);
       }
 
-      RAJA_FT_END;
     }
 
     return run_storage;
diff --git a/include/RAJA/policy/hip/forall.hpp b/include/RAJA/policy/hip/forall.hpp
@@ -552,7 +552,6 @@ forall_impl(resources::Hip hip_res,
     internal::HipDims dims(1);
     DimensionCalculator::set_dimensions(dims, len, func, shmem);
 
-    RAJA_FT_BEGIN;
 
     RAJA::hip::detail::hipInfo launch_info;
     launch_info.gridDim  = dims.blocks;
@@ -581,7 +580,6 @@ forall_impl(resources::Hip hip_res,
                                                       launch_info);
     }
 
-    RAJA_FT_END;
   }
 
   return resources::EventProxy<resources::Hip>(hip_res);
diff --git a/include/RAJA/policy/hip/launch.hpp b/include/RAJA/policy/hip/launch.hpp
@@ -93,7 +93,6 @@ struct LaunchExecute<
         blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       size_t shared_mem_size = launch_params.shared_mem_size;
       RAJA::hip::detail::hipInfo launch_info;
@@ -125,7 +124,6 @@ struct LaunchExecute<
                                                         launch_info);
       }
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
@@ -200,7 +198,6 @@ struct LaunchExecute<RAJA::policy::hip::hip_launch_t<async, nthreads>>
         blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       size_t shared_mem_size = launch_params.shared_mem_size;
       RAJA::hip::detail::hipInfo launch_info;
@@ -232,7 +229,6 @@ struct LaunchExecute<RAJA::policy::hip::hip_launch_t<async, nthreads>>
                                                         launch_info);
       }
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
diff --git a/include/RAJA/policy/sycl/launch.hpp b/include/RAJA/policy/sycl/launch.hpp
@@ -70,7 +70,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         params.teams.value[1] > zero && params.teams.value[2] > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       q->submit([&](::sycl::handler& h) {
         auto s_vec = ::sycl::local_accessor<char, 1>(params.shared_mem_size, h);
@@ -94,7 +93,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         q->wait();
       }
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
@@ -152,7 +150,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         return x;
       };
 
-      RAJA_FT_BEGIN;
 
       ReduceParams* res = ::sycl::malloc_shared<ReduceParams>(1, *q);
       RAJA::expt::ParamMultiplexer::parampack_init(pol, *res);
@@ -185,7 +182,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
                                                       *res);
       ::sycl::free(res, *q);
 
-      RAJA_FT_END;
     }
 
     RAJA::expt::ParamMultiplexer::parampack_resolve(pol, launch_reducers);
@@ -231,7 +227,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         params.teams.value[1] > zero && params.teams.value[2] > zero)
     {
 
-      RAJA_FT_BEGIN;
 
       //
       // Kernel body is nontrivially copyable, create space on device and copy
@@ -262,7 +257,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
 
       ::sycl::free(lbody, *q);
 
-      RAJA_FT_END;
     }
 
     return resources::EventProxy<resources::Resource>(res);
@@ -319,7 +313,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
         return x;
       };
 
-      RAJA_FT_BEGIN;
 
       //
       // Kernel body is nontrivially copyable, create space on device and copy
@@ -362,7 +355,6 @@ struct LaunchExecute<RAJA::sycl_launch_t<async, 0>>
       ::sycl::free(res, *q);
       ::sycl::free(lbody, *q);
 
-      RAJA_FT_END;
     }
 
     RAJA::expt::ParamMultiplexer::parampack_resolve(pol, launch_reducers);

Original file line number	Diff line number	Diff line change
`@@ -377,7 +377,6 @@ struct WorkRunner<`
`377`	`377`	`block_size),`
`378`	`378`	`static_cast<cuda_dim_member_t>(num_loops), 1};`
`379`	`379`
`380`		`- RAJA_FT_BEGIN;`
`381`	`380`
`382`	`381`	`//`
`383`	`382`	`// Setup shared memory buffers`
`@@ -393,7 +392,6 @@ struct WorkRunner<`
`393`	`392`	`shmem, r, Async);`
`394`	`393`	`}`
`395`	`394`
`396`		`- RAJA_FT_END;`
`397`	`395`	`}`
`398`	`396`
`399`	`397`	`return run_storage;`
Original file line number	Diff line number	Diff line change
`@@ -570,7 +570,6 @@ forall_impl(resources::Cuda cuda_res,`
`570`	`570`	`internal::CudaDims dims(1);`
`571`	`571`	`DimensionCalculator::set_dimensions(dims, len, func, shmem);`
`572`	`572`
`573`		`- RAJA_FT_BEGIN;`
`574`	`573`
`575`	`574`	`RAJA::cuda::detail::cudaInfo launch_info;`
`576`	`575`	`launch_info.gridDim = dims.blocks;`
`@@ -599,7 +598,6 @@ forall_impl(resources::Cuda cuda_res,`
`599`	`598`	`launch_info);`
`600`	`599`	`}`
`601`	`600`
`602`		`- RAJA_FT_END;`
`603`	`601`	`}`
`604`	`602`
`605`	`603`	`return resources::EventProxy<resources::Cuda>(cuda_res);`
Original file line number	Diff line number	Diff line change
`@@ -96,7 +96,6 @@ struct LaunchExecute<`
`96`	`96`	`blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)`
`97`	`97`	`{`
`98`	`98`
`99`		`- RAJA_FT_BEGIN;`
`100`	`99`
`101`	`100`	`size_t shared_mem_size = launch_params.shared_mem_size;`
`102`	`101`	`RAJA::cuda::detail::cudaInfo launch_info;`
`@@ -127,7 +126,6 @@ struct LaunchExecute<`
`127`	`126`	`launch_info);`
`128`	`127`	`}`
`129`	`128`
`130`		`- RAJA_FT_END;`
`131`	`129`	`}`
`132`	`130`
`133`	`131`	`return resources::EventProxy<resources::Resource>(res);`
`@@ -207,7 +205,6 @@ struct LaunchExecute<`
`207`	`205`	`blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)`
`208`	`206`	`{`
`209`	`207`
`210`		`- RAJA_FT_BEGIN;`
`211`	`208`
`212`	`209`	`size_t shared_mem_size = launch_params.shared_mem_size;`
`213`	`210`	`RAJA::cuda::detail::cudaInfo launch_info;`
`@@ -238,7 +235,6 @@ struct LaunchExecute<`
`238`	`235`	`launch_info);`
`239`	`236`	`}`
`240`	`237`
`241`		`- RAJA_FT_END;`
`242`	`238`	`}`
`243`	`239`
`244`	`240`	`return resources::EventProxy<resources::Resource>(res);`
Original file line number	Diff line number	Diff line change
`@@ -364,7 +364,6 @@ struct WorkRunner<`
`364`	`364`	`block_size),`
`365`	`365`	`static_cast<hip_dim_member_t>(num_loops), 1};`
`366`	`366`
`367`		`- RAJA_FT_BEGIN;`
`368`	`367`
`369`	`368`	`//`
`370`	`369`	`// Setup shared memory buffers`
`@@ -380,7 +379,6 @@ struct WorkRunner<`
`380`	`379`	`shmem, r, Async);`
`381`	`380`	`}`
`382`	`381`
`383`		`- RAJA_FT_END;`
`384`	`382`	`}`
`385`	`383`
`386`	`384`	`return run_storage;`
Original file line number	Diff line number	Diff line change
`@@ -552,7 +552,6 @@ forall_impl(resources::Hip hip_res,`
`552`	`552`	`internal::HipDims dims(1);`
`553`	`553`	`DimensionCalculator::set_dimensions(dims, len, func, shmem);`
`554`	`554`
`555`		`- RAJA_FT_BEGIN;`
`556`	`555`
`557`	`556`	`RAJA::hip::detail::hipInfo launch_info;`
`558`	`557`	`launch_info.gridDim = dims.blocks;`
`@@ -581,7 +580,6 @@ forall_impl(resources::Hip hip_res,`
`581`	`580`	`launch_info);`
`582`	`581`	`}`
`583`	`582`
`584`		`- RAJA_FT_END;`
`585`	`583`	`}`
`586`	`584`
`587`	`585`	`return resources::EventProxy<resources::Hip>(hip_res);`
Original file line number	Diff line number	Diff line change
`@@ -93,7 +93,6 @@ struct LaunchExecute<`
`93`	`93`	`blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)`
`94`	`94`	`{`
`95`	`95`
`96`		`- RAJA_FT_BEGIN;`
`97`	`96`
`98`	`97`	`size_t shared_mem_size = launch_params.shared_mem_size;`
`99`	`98`	`RAJA::hip::detail::hipInfo launch_info;`
`@@ -125,7 +124,6 @@ struct LaunchExecute<`
`125`	`124`	`launch_info);`
`126`	`125`	`}`
`127`	`126`
`128`		`- RAJA_FT_END;`
`129`	`127`	`}`
`130`	`128`
`131`	`129`	`return resources::EventProxy<resources::Resource>(res);`
`@@ -200,7 +198,6 @@ struct LaunchExecute<RAJA::policy::hip::hip_launch_t<async, nthreads>>`
`200`	`198`	`blockSize.x > zero && blockSize.y > zero && blockSize.z > zero)`
`201`	`199`	`{`
`202`	`200`
`203`		`- RAJA_FT_BEGIN;`
`204`	`201`
`205`	`202`	`size_t shared_mem_size = launch_params.shared_mem_size;`
`206`	`203`	`RAJA::hip::detail::hipInfo launch_info;`
`@@ -232,7 +229,6 @@ struct LaunchExecute<RAJA::policy::hip::hip_launch_t<async, nthreads>>`
`232`	`229`	`launch_info);`
`233`	`230`	`}`
`234`	`231`
`235`		`- RAJA_FT_END;`
`236`	`232`	`}`
`237`	`233`
`238`	`234`	`return resources::EventProxy<resources::Resource>(res);`