Hardcode input/output dimensions of local laplacian

Author: antonysigma

apps/harris/CMakeLists.txt

@@ -16,7 +16,7 @@ add_halide_generator(harris.generator SOURCES harris_generator.cpp)
 
 set(_harris_autoscheduler_params autoscheduler.experimental_gpu_schedule=1)
 
-if(Halide_TARGET MATCHES "opencl|metal")
+if(Halide_TARGET MATCHES "opencl|metal|cuda|vulkan")
     # Set last_level_cache per GPU block to an extremely small value. This
     # eliminates all `.compute_at` in the generated schedules, which in turn
     # eliminates all GPU shared memory allocations.

apps/local_laplacian/local_laplacian_generator.cpp

@@ -97,6 +97,19 @@ class LocalLaplacian : public Halide::Generator<LocalLaplacian> {
         // Provide estimates on the pipeline output
         output.set_estimates({{0, 1536}, {0, 2560}, {0, 3}});
 
+        // Hardcode the input and output dimensions to suppress the OpenCL/Metal
+        // launch failure:
+        //
+        //  OpenCL error: CL_INVALID_WORK_GROUP_SIZE clEnqueueNDRangeKernel
+        //  failed
+        input.dim(0).set_bounds(0, 1536).set_stride(1);
+        input.dim(1).set_bounds(0, 2560).set_stride(1536);
+        input.dim(2).set_bounds(0, 3).set_stride(1536 * 2560);
+
+        output.dim(0).set_bounds(0, 1536).set_stride(1);
+        output.dim(1).set_bounds(0, 2560).set_stride(1536);
+        output.dim(2).set_bounds(0, 3).set_stride(1536 * 2560);
+
         /* THE SCHEDULE */
         if (using_autoscheduler()) {
             // Nothing.