From 9834c07c4294a018d7b0df5c00c1054e585dfcd0 Mon Sep 17 00:00:00 2001
From: Sven Verdoolaege <skimo@kotnet.org>
Date: Wed, 12 Jun 2013 09:59:50 +0200
Subject: [PATCH] ppcg_kernel: remove fields referring to original grid size

Since the kernel launch has been modified to take into account
the effective grid size, the original grid size is not relevant
inside the kernel.

Signed-off-by: Sven Verdoolaege <skimo@kotnet.org>
---
 cuda.c | 9 +++++----
 gpu.c  | 3 ---
 gpu.h  | 3 ---
 3 files changed, 5 insertions(+), 10 deletions(-)

diff --git a/cuda.c b/cuda.c
index 0797eeb..18ec802 100644
--- a/cuda.c
+++ b/cuda.c
@@ -311,7 +311,7 @@ static void print_indent(FILE *dst, int indent)
 
 static void print_kernel_iterators(FILE *out, struct ppcg_kernel *kernel)
 {
-	int i;
+	int i, n_grid;
 	isl_ctx *ctx = isl_ast_node_get_ctx(kernel->tree);
 	const char *type;
 	const char *block_dims[] = { "blockIdx.x", "blockIdx.y" };
@@ -320,14 +320,15 @@ static void print_kernel_iterators(FILE *out, struct ppcg_kernel *kernel)
 
 	type = isl_options_get_ast_iterator_type(ctx);
 
-	if (kernel->n_grid > 0) {
+	n_grid = isl_multi_pw_aff_dim(kernel->grid_size, isl_dim_set);
+	if (n_grid > 0) {
 		print_indent(out, 4);
 		fprintf(out, "%s ", type);
-		for (i = 0; i < kernel->n_grid; ++i) {
+		for (i = 0; i < n_grid; ++i) {
 			if (i)
 				fprintf(out, ", ");
 			fprintf(out, "b%d = %s",
-				i, block_dims[kernel->n_grid - 1 - i]);
+				i, block_dims[n_grid - 1 - i]);
 		}
 		fprintf(out, ";\n");
 	}
diff --git a/gpu.c b/gpu.c
index cbfb115..d0f93eb 100644
--- a/gpu.c
+++ b/gpu.c
@@ -4303,9 +4303,6 @@ static __isl_give isl_ast_node *create_host_leaf(
 	kernel->n_block = gen->n_block;
 	for (i = 0; i < gen->n_block; ++i)
 		kernel->block_dim[i] = gen->block_dim[i];
-	kernel->n_grid = gen->n_grid;
-	for (i = 0; i < gen->n_grid; ++i)
-		kernel->grid_dim[i] = gen->grid_dim[i];
 	kernel->context = isl_union_map_params(isl_union_map_copy(schedule));
 	kernel->grid_size = extract_grid_size(gen, kernel);
 	kernel->arrays = isl_union_map_range(access);
diff --git a/gpu.h b/gpu.h
index c891286..0bd6f50 100644
--- a/gpu.h
+++ b/gpu.h
@@ -180,9 +180,6 @@ struct ppcg_kernel_var {
 struct ppcg_kernel {
 	int id;
 
-	int n_grid;
-	int grid_dim[2];
-
 	int n_block;
 	int block_dim[3];
 
-- 
2.11.4.GIT