Review: add brackets for readability, remove ggml_set_param and add asserts

Author: am17an

ggml/src/ggml-cuda/conv2d-transpose.cu

@@ -16,8 +16,8 @@ __global__ void conv2d_transpose_kernel(const float * __restrict__ input, const
     }
 
     const int out_x_idx = global_idx % out_w;
-    const int out_y_idx = global_idx / out_w % out_h;
-    const int c_idx     = global_idx / (out_w * out_h) % c_out;
+    const int out_y_idx = (global_idx / out_w) % out_h;
+    const int c_idx     = (global_idx / (out_w * out_h)) % c_out;
     const int n_idx     = global_idx / (out_w * out_h * c_out);
 
     float accumulator = 0;
@@ -78,6 +78,10 @@ void ggml_cuda_conv_2d_transpose_p0(ggml_backend_cuda_context & ctx, ggml_tensor
 
     cudaStream_t st = ctx.stream();
 
+    GGML_ASSERT(ggml_is_contiguous(input));
+    GGML_ASSERT(ggml_is_contiguous(kernel));
+    GGML_ASSERT(ggml_is_contiguous(dst));
+
     const int total  = (output_w * output_h * channels_out * batches);
     const int blocks = (total + CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE - 1) / CUDA_CONV2D_TRANSPOSE_BLOCK_SIZE;
 

tests/test-backend-ops.cpp

@@ -560,7 +560,6 @@ struct test_case {
             }
 
             double err = nmse(f1.data(), f2.data(), f1.size());
-
             if (err > ud->max_err) {
                 printf("[%s] NMSE = %.9f > %.9f ", ggml_op_desc(t1), err, ud->max_err);
                 //for (int i = 0; i < (int) f1.size(); i++) {
@@ -2743,11 +2742,9 @@ struct test_conv_transpose_2d : public test_case {
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         ggml_tensor * input = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne_input.data());
-        ggml_set_param(input);
         ggml_set_name(input, "input");
 
         ggml_tensor * kernel = ggml_new_tensor(ctx, GGML_TYPE_F16, 4, ne_kernel.data());
-        ggml_set_param(kernel);
         ggml_set_name(kernel, "kernel");
 
         ggml_tensor * out = ggml_conv_transpose_2d_p0(ctx, kernel, input, stride);