[GPGPU] Simplify PPCGSCop to reduce compile time [NFC]

[grosser]

Summary

Drop unused parameter dimensions to reduce the size of the sets we are working
with. Especially the computed dependences tend to accumulate a lot of parameters
that are present in the input memory accesses, but often not necessary to
express the actual dependences. As isl represents maps and sets with dense
matrices, reducing the dimensionality of isl sets commonly reduces code
generation performance.

This reduces compile time from 17 to 11 seconds for our test case. While this is
not impressive, this patch helped me to identify the previous two performance
improvements and additionally also increases readability of the isl data
structures we use.

[bollu]

Other than comments, LGTM :)

lib/CodeGen/PPCGCodeGeneration.cpp

2639 ↗

(On Diff #111651)

Please name this with the function intent. suggestion: RemoveUnusedDims.

2643 ↗

(On Diff #111651)

const int Dim = i - Removed.

2653 ↗

(On Diff #111651)

Alternate way without having to name the Lambda, just move the lambda into the callsite?

UMap.foreach_map([&New](...) {...});

2661 ↗

(On Diff #111651)

Similar comment: RemoveUnusedDims, maybe.

2665 ↗

(On Diff #111651)

const int Dim = i - Removed.

2675 ↗

(On Diff #111651)

Similar comment: Consider moving Lambda into callsite?

3207 ↗

(On Diff #111651)

Nit: Move this about the has_permutable?

polly/trunk/lib/CodeGen/PPCGCodeGeneration.cpp

for (auto &Array : S->arrays()) {
       Names = isl_id_to_ast_expr_set(Names, Id, isl_ast_expr_copy(Zero));
     }
 
     isl_ast_expr_free(Zero);
 
     return Names;
   }
 
+  /// Remove unreferenced parameter dimensions from union_map.
+  isl::union_map removeUnusedParameters(isl::union_map UMap) {
+    auto New = isl::union_map::empty(isl::space(UMap.get_ctx(), 0, 0));
+
+    auto RemoveUnusedDims = [&New](isl::map S) -> isl::stat {
+      int Removed = 0;
+      int NumDims = S.dim(isl::dim::param);
+      for (long i = 0; i < NumDims; i++) {
+        const int Dim = i - Removed;
+        if (!S.involves_dims(isl::dim::param, Dim, 1)) {
+          S = S.remove_dims(isl::dim::param, Dim, 1);
+          Removed++;
+        }
+      }
+      New = New.unite(S);
+      return isl::stat::ok;
+    };
+
+    UMap.foreach_map(RemoveUnusedDims);
+    return New;
+  }
+
+  /// Remove unreferenced parameter dimensions from union_set.
+  isl::union_set removeUnusedParameters(isl::union_set USet) {
+    auto New = isl::union_set::empty(isl::space(USet.get_ctx(), 0, 0));
+
+    auto RemoveUnusedDims = [&New](isl::set S) -> isl::stat {
+      int Removed = 0;
+      int NumDims = S.dim(isl::dim::param);
+      for (long i = 0; i < NumDims; i++) {
+        const int Dim = i - Removed;
+        if (!S.involves_dims(isl::dim::param, Dim, 1)) {
+          S = S.remove_dims(isl::dim::param, Dim, 1);
+          Removed++;
+        }
+      }
+      New = New.unite(S);
+      return isl::stat::ok;
+    };
+
+    USet.foreach_set(RemoveUnusedDims);
+    return New;
+  }
+
+  /// Simplify PPCG scop to improve compile time.
+  ///
+  /// We drop unused parameter dimensions to reduce the size of the sets we are
+  /// working with. Especially the computed dependences tend to accumulate a lot
+  /// of parameters that are present in the input memory accesses, but often are
+  /// not necessary to express the actual dependences. As isl represents maps
+  /// and sets with dense matrices, reducing the dimensionality of isl sets
+  /// commonly reduces code generation performance.
+  void simplifyPPCGScop(ppcg_scop *PPCGScop) {
+    PPCGScop->domain =
+        removeUnusedParameters(isl::manage(PPCGScop->domain)).release();
+
+    PPCGScop->dep_forced =
+        removeUnusedParameters(isl::manage(PPCGScop->dep_forced)).release();
+    PPCGScop->dep_false =
+        removeUnusedParameters(isl::manage(PPCGScop->dep_false)).release();
+    PPCGScop->dep_flow =
+        removeUnusedParameters(isl::manage(PPCGScop->dep_flow)).release();
+    PPCGScop->tagged_dep_flow =
+        removeUnusedParameters(isl::manage(PPCGScop->tagged_dep_flow))
+            .release();
+
+    PPCGScop->tagged_dep_order =
+        removeUnusedParameters(isl::manage(PPCGScop->tagged_dep_order))
+            .release();
+  }
+
   /// Create a new PPCG scop from the current scop.
   ///
   /// The PPCG scop is initialized with data from the current polly::Scop. From
   /// this initial data, the data-dependences in the PPCG scop are initialized.
   /// We do not use Polly's dependence analysis for now, to ensure we match
   /// the PPCG default behaviour more closely.
   ///
   /// @returns A new ppcg scop.
if (KillsInfo.KillsSchedule.get())
           PPCGScop->schedule, KillsInfo.KillsSchedule.take());
 
     PPCGScop->names = getNames();
     PPCGScop->pet = nullptr;
 
     compute_tagger(PPCGScop);
     compute_dependences(PPCGScop);
     eliminate_dead_code(PPCGScop);
+    simplifyPPCGScop(PPCGScop);
 
     return PPCGScop;
   }
 
   /// Collect the array accesses in a statement.
   ///
   /// @param Stmt The statement for which to collect the accesses.
   ///
gpu_gen *generateGPU(ppcg_scop *PPCGScop, gpu_prog *PPCGProg) {
 
     // Set scheduling strategy to same strategy PPCG is using.
     isl_options_set_schedule_outer_coincidence(PPCGGen->ctx, true);
     isl_options_set_schedule_maximize_band_depth(PPCGGen->ctx, true);
     isl_options_set_schedule_whole_component(PPCGGen->ctx, false);
 
     isl_schedule *Schedule = get_schedule(PPCGGen);
 
-    int has_permutable = has_any_permutable_node(Schedule);
-
+    /// Copy to and from device functions may introduce new parameters, which
+    /// must be present in the schedule tree root for code generation. Hence,
+    /// we ensure that all possible parameters are introduced from this point.
+    if (!PollyManagedMemory)
       Schedule =
           isl_schedule_align_params(Schedule, S->getFullParamSpace().release());
 
+    int has_permutable = has_any_permutable_node(Schedule);
+
     if (!has_permutable || has_permutable < 0) {
       Schedule = isl_schedule_free(Schedule);
       DEBUG(dbgs() << getUniqueScopName(S)
                    << " does not have permutable bands. Bailing out\n";);
     } else {
       const bool CreateTransferToFromDevice = !PollyManagedMemory;
       Schedule = map_to_device(PPCGGen, Schedule, CreateTransferToFromDevice);
       PPCGGen->tree = generate_code(PPCGGen, isl_schedule_copy(Schedule));

polly/trunk/test/GPGPU/kernel-params-only-some-arrays.ll

 ;        B[i] += 42;
 ;    }
 
 ; KERNEL: ; ModuleID = 'FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_0'
 ; KERNEL-NEXT: source_filename = "FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_0"
 ; KERNEL-NEXT: target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 ; KERNEL-NEXT: target triple = "nvptx64-nvidia-cuda"
 
-; KERNEL: define ptx_kernel void @FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_0(i8 addrspace(1)* %MemRef_B)
+; KERNEL: define ptx_kernel void @FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_0(i8 addrspace(1)* %MemRef_A)
 ; KERNEL-NEXT:   entry:
 ; KERNEL-NEXT:     %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 ; KERNEL-NEXT:     %b0 = zext i32 %0 to i64
 ; KERNEL-NEXT:     %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 ; KERNEL-NEXT:     %t0 = zext i32 %1 to i64
 
 ; KERNEL:     ret void
 ; KERNEL-NEXT: }
 
 ; KERNEL: ; ModuleID = 'FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_1'
 ; KERNEL-NEXT: source_filename = "FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_1"
 ; KERNEL-NEXT: target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-i128:128:128-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
 ; KERNEL-NEXT: target triple = "nvptx64-nvidia-cuda"
 
-; KERNEL: define ptx_kernel void @FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_1(i8 addrspace(1)* %MemRef_A)
+; KERNEL: define ptx_kernel void @FUNC_kernel_params_only_some_arrays_SCOP_0_KERNEL_1(i8 addrspace(1)* %MemRef_B)
 ; KERNEL-NEXT:   entry:
 ; KERNEL-NEXT:     %0 = call i32 @llvm.nvvm.read.ptx.sreg.ctaid.x()
 ; KERNEL-NEXT:     %b0 = zext i32 %0 to i64
 ; KERNEL-NEXT:     %1 = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
 ; KERNEL-NEXT:     %t0 = zext i32 %1 to i64
 
 ; KERNEL:     ret void
 ; KERNEL-NEXT: }
 
 
-; IR:       [[DEVPTR:%.*]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_B)
+; IR:       [[DEVPTR:%.*]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_A)
 ; IR-NEXT:  [[SLOT:%.*]] = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_0_params, i64 0, i64 0
 ; IR-NEXT:  store i8* [[DEVPTR]], i8** %polly_launch_0_param_0
 ; IR-NEXT:  [[DATA:%.*]] = bitcast i8** %polly_launch_0_param_0 to i8*
 ; IR-NEXT:  store i8* [[DATA]], i8** [[SLOT]]
 
-; IR:       [[DEVPTR:%.*]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_A)
+; IR:       [[DEVPTR:%.*]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_B)
 ; IR-NEXT:  [[SLOT:%.*]] = getelementptr [2 x i8*], [2 x i8*]* %polly_launch_1_params, i64 0, i64 0
 ; IR-NEXT:  store i8* [[DEVPTR]], i8** %polly_launch_1_param_0
 ; IR-NEXT:  [[DATA:%.*]] = bitcast i8** %polly_launch_1_param_0 to i8*
 ; IR-NEXT:  store i8* [[DATA]], i8** [[SLOT]]
 
 
 target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128"
 

polly/trunk/test/GPGPU/non-zero-array-offset.ll

 ; RUN: opt %loadPolly -polly-codegen-ppcg -polly-acc-dump-code \
 ; RUN: -disable-output < %s | \
 ; RUN: FileCheck -check-prefix=CODE %s
 
 ; RUN: opt %loadPolly -polly-codegen-ppcg -S < %s | \
 ; RUN: FileCheck %s -check-prefix=IR
 ;
 ; REQUIRES: pollyacc
 
 ; CODE:      cudaCheckReturn(cudaMemcpy(dev_MemRef_B, MemRef_B, (16) * sizeof(float), cudaMemcpyHostToDevice));
 ; CODE-NEXT: cudaCheckReturn(cudaMemcpy(dev_MemRef_A, MemRef_A, (8) * sizeof(float), cudaMemcpyHostToDevice));
 
 ; CODE:          dim3 k0_dimBlock(8);
 ; CODE-NEXT:     dim3 k0_dimGrid(1);
-; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_A);
+; CODE-NEXT:     kernel0 <<<k0_dimGrid, k0_dimBlock>>> (dev_MemRef_B);
 ; CODE-NEXT:     cudaCheckKernel();
 ; CODE-NEXT:   }
 
 ; CODE:        {
 ; CODE-NEXT:     dim3 k1_dimBlock(8);
 ; CODE-NEXT:     dim3 k1_dimGrid(1);
-; CODE-NEXT:     kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_B);
+; CODE-NEXT:     kernel1 <<<k1_dimGrid, k1_dimBlock>>> (dev_MemRef_A);
 ; CODE-NEXT:     cudaCheckKernel();
 ; CODE-NEXT:   }
 
 ; CODE:   cudaCheckReturn(cudaMemcpy(MemRef_B, dev_MemRef_B, (16) * sizeof(float), cudaMemcpyDeviceToHost));
 ; CODE-NEXT:   cudaCheckReturn(cudaMemcpy(MemRef_A, dev_MemRef_A, (8) * sizeof(float), cudaMemcpyDeviceToHost));
 
 ; CODE: # kernel0
-; CODE-NEXT: Stmt_bb11(t0);
+; CODE-NEXT: Stmt_bb3(t0);
 
 ; CODE: # kernel1
-; CODE-NEXT: Stmt_bb3(t0);
+; CODE-NEXT: Stmt_bb11(t0);
 
 ; IR:       %p_dev_array_MemRef_B = call i8* @polly_allocateMemoryForDevice(i64 32)
 ; IR-NEXT:  %p_dev_array_MemRef_A = call i8* @polly_allocateMemoryForDevice(i64 32)
 ; IR-NEXT:  [[REG0:%.+]] = getelementptr float, float* %B, i64 8
 ; IR-NEXT:  [[REG1:%.+]] = bitcast float* [[REG0]] to i8*
 ; IR-NEXT:  call void @polly_copyFromHostToDevice(i8* [[REG1]], i8* %p_dev_array_MemRef_B, i64 32)
 
 ; IR:      [[REGA:%.+]] = call i8* @polly_getDevicePtr(i8* %p_dev_array_MemRef_B)