Update compiletest to edition 2024.

Again, all the unsafe operations within unsafe functions now need
explicit `unsafe` marking.

Author: nnethercote

tests/compiletests/src/main.rs

@@ -109,7 +109,7 @@ impl Runner {
                     .iter()
                     .map(|dir| format!("-L dependency={}", dir.display()))
                     .fold(String::new(), |a, b| b + " " + &a),
-                "--edition 2021",
+                "--edition 2024",
                 &*format!("--extern noprelude:core={}", deps.core.display()),
                 &*format!(
                     "--extern noprelude:compiler_builtins={}",

tests/compiletests/ui/atomic/atomic_operations.rs

@@ -17,7 +17,7 @@ pub unsafe fn test_cuda_atomic_floats() {
 
     // Block-scoped atomic float
     let block_atomic = BlockAtomicF32::new(1.5);
-    let _old = block_atomic.fetch_add(0.5, Ordering::Relaxed);
+    let _old = unsafe { block_atomic.fetch_add(0.5, Ordering::Relaxed) };
 
     // System-scoped atomic float
     let system_atomic = SystemAtomicF32::new(0.0);
@@ -29,7 +29,7 @@ pub unsafe fn test_cuda_atomic_floats() {
 
     // Test block-scoped f64
     let block_f64 = BlockAtomicF64::new(2.718);
-    let _old = block_f64.fetch_sub(0.5, Ordering::Relaxed);
+    let _old = unsafe { block_f64.fetch_sub(0.5, Ordering::Relaxed) };
 
     // Test bitwise operations on floats
     let _old = atomic_f32.fetch_and(3.14, Ordering::Relaxed);

tests/compiletests/ui/core/ops/range_contains.rs

@@ -8,5 +8,7 @@ fn has_two_decimal_digits(x: u32) -> bool {
 
 #[kernel]
 pub unsafe fn main(i: u32, o: *mut u32) {
-    *o = has_two_decimal_digits(i) as u32;
+    unsafe {
+        *o = has_two_decimal_digits(i) as u32;
+    }
 }

tests/compiletests/ui/dis/simple_add.rs

@@ -7,6 +7,8 @@ use cuda_std::kernel;
 
 #[kernel]
 pub unsafe fn simple_add_kernel(a: *const f32, b: *const f32, c: *mut f32) {
-    let sum = *a + *b;
-    *c = sum;
+    unsafe {
+        let sum = *a + *b;
+        *c = sum;
+    }
 }

tests/compiletests/ui/dis/simple_add.stderr

@@ -19,17 +19,17 @@ $L__tmp0:
     cvta.to.global.u64     %rd4, %rd3;
     cvta.to.global.u64     %rd5, %rd2;
     cvta.to.global.u64     %rd6, %rd1;
-    .loc    1 10 15
+    .loc    1 11 19
     ld.global.f32     %f1, [%rd6];
-    .loc    1 10 20
+    .loc    1 11 24
     ld.global.f32     %f2, [%rd5];
-    .loc    1 10 15
+    .loc    1 11 19
     add.f32     %f3, %f1, %f2;
 $L__tmp1:
-    .loc    1 11 5
+    .loc    1 12 9
     st.global.f32     [%rd4], %f3;
 $L__tmp2:
-    .loc    1 12 2
+    .loc    1 14 2
     ret;
 $L__tmp3:
 $L__func_end0:

tests/compiletests/ui/glam/mat3_vec3_multiply.rs

@@ -7,5 +7,7 @@ use cuda_std::kernel;
 #[kernel]
 pub unsafe fn mat3_vec3_multiply(input: glam::Mat3, output: *mut glam::Vec3) {
     let vector = input * glam::Vec3::new(1.0, 2.0, 3.0);
-    *output = vector;
+    unsafe {
+        *output = vector;
+    }
 }

tests/compiletests/ui/glam/mat4_operations.rs

@@ -16,15 +16,21 @@ pub unsafe fn mat4_transform_operations(
 ) {
     // Transform a 3D point (w=1 implied)
     let transformed_point = matrix.transform_point3(point);
-    *result_point = transformed_point;
+    unsafe {
+        *result_point = transformed_point;
+    }
 
     // Transform a 4D vector
     let transformed_vector = matrix * vector;
-    *result_vector = transformed_vector;
+    unsafe {
+        *result_vector = transformed_vector;
+    }
 
     // Calculate determinant
     let det = matrix.determinant();
-    *result_determinant = det;
+    unsafe {
+        *result_determinant = det;
+    }
 }
 
 #[kernel]
@@ -39,13 +45,19 @@ pub unsafe fn mat4_construction(
 ) {
     // Create translation matrix
     let trans_mat = Mat4::from_translation(translation);
-    *result_translation = trans_mat;
+    unsafe {
+        *result_translation = trans_mat;
+    }
 
     // Create scale matrix
     let scale_mat = Mat4::from_scale(scale);
-    *result_scale = scale_mat;
+    unsafe {
+        *result_scale = scale_mat;
+    }
 
     // Create rotation matrix
     let rot_mat = Mat4::from_axis_angle(axis, angle_radians);
-    *result_rotation = rot_mat;
+    unsafe {
+        *result_rotation = rot_mat;
+    }
 }

tests/compiletests/ui/glam/vec3_operations.rs

@@ -15,15 +15,21 @@ pub unsafe fn vec3_basic_ops(
 ) {
     // Vector addition
     let sum = a + b;
-    *result_add = sum;
+    unsafe {
+        *result_add = sum;
+    }
 
     // Dot product
     let dot = a.dot(b);
-    *result_dot = dot;
+    unsafe {
+        *result_dot = dot;
+    }
 
     // Cross product
     let cross = a.cross(b);
-    *result_cross = cross;
+    unsafe {
+        *result_cross = cross;
+    }
 }
 
 #[kernel]
@@ -34,9 +40,13 @@ pub unsafe fn vec3_normalization(
 ) {
     // Get length
     let len = input.length();
-    *result_length = len;
+    unsafe {
+        *result_length = len;
+    }
 
     // Normalize
     let normalized = input.normalize();
-    *result_normalized = normalized;
+    unsafe {
+        *result_normalized = normalized;
+    }
 }

tests/compiletests/ui/hello_world.rs

@@ -4,5 +4,7 @@ use cuda_std::kernel;
 
 #[kernel]
 pub unsafe fn add_one(x: *mut f32) {
-    *x = *x + 1.0;
+    unsafe {
+        *x = *x + 1.0;
+    }
 }

tests/compiletests/ui/lang/consts/constant_memory_overflow.rs

@@ -17,5 +17,5 @@ static BIG_ARRAY_3: [u32; ARRAY_SIZE] = [333u32; ARRAY_SIZE];
 
 #[kernel]
 pub unsafe fn test_kernel(out: *mut u32) {
-    *out = BIG_ARRAY_1[0] + BIG_ARRAY_2[0] + BIG_ARRAY_3[0];
+    unsafe { *out = BIG_ARRAY_1[0] + BIG_ARRAY_2[0] + BIG_ARRAY_3[0] };
 }