Skip to content

Commit eb737a7

Browse files
author
Jessica Shi
committed
fix miscellaneous issues
1 parent 1e7c6c5 commit eb737a7

File tree

9 files changed

+101
-1321
lines changed

9 files changed

+101
-1321
lines changed

‎examples/spmv_assembly.c

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// Generated by the Tensor Algebra Compiler (tensor-compiler.org)
2-
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
2+
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -s=split(i,i0,i1,32) -s=reorder(i0,i1,j) -s=parallelize(i0,CPUThread,NoRaces) -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
33

44
int assemble(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
55
int y1_dimension = (int)(y->dimensions[0]);
@@ -9,4 +9,4 @@ int assemble(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
99

1010
y->vals = (uint8_t*)y_vals;
1111
return 0;
12-
}
12+
}

‎examples/spmv_compute.c

Lines changed: 17 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// Generated by the Tensor Algebra Compiler (tensor-compiler.org)
2-
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
2+
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -s=split(i,i0,i1,32) -s=reorder(i0,i1,j) -s=parallelize(i0,CPUThread,NoRaces) -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
33

44
int compute(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
55
int y1_dimension = (int)(y->dimensions[0]);
@@ -11,14 +11,23 @@ int compute(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
1111
int x1_dimension = (int)(x->dimensions[0]);
1212
double* restrict x_vals = (double*)(x->vals);
1313

14+
#pragma omp parallel for schedule(static)
15+
for (int32_t py = 0; py < y1_dimension; py++) {
16+
y_vals[py] = 0.0;
17+
}
18+
1419
#pragma omp parallel for schedule(runtime)
15-
for (int32_t i = 0; i < A1_dimension; i++) {
16-
double y_val = 0.0;
17-
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
18-
int32_t j = A2_crd[jA];
19-
y_val += A_vals[jA] * x_vals[j];
20+
for (int32_t i0 = 0; i0 < ((A1_dimension + 31) / 32); i0++) {
21+
for (int32_t i1 = 0; i1 < 32; i1++) {
22+
int32_t i = i0 * 32 + i1;
23+
if (i >= A1_dimension)
24+
continue;
25+
26+
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
27+
int32_t j = A2_crd[jA];
28+
y_vals[i] = y_vals[i] + A_vals[jA] * x_vals[j];
29+
}
2030
}
21-
y_vals[i] = y_val;
2231
}
2332
return 0;
24-
}
33+
}

‎examples/spmv_full.c

Lines changed: 35 additions & 17 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
// Generated by the Tensor Algebra Compiler (tensor-compiler.org)
2-
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
2+
// taco "y(i)=A(i,j)*x(j)" -f=y:d:0 -f=A:ds:0,1 -f=x:d:0 -s=split(i,i0,i1,32) -s=reorder(i0,i1,j) -s=parallelize(i0,CPUThread,NoRaces) -write-source=taco_kernel.c -write-compute=taco_compute.c -write-assembly=taco_assembly.c
33
#ifndef TACO_C_HEADERS
44
#define TACO_C_HEADERS
55
#include <stdio.h>
@@ -118,14 +118,23 @@ int compute(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
118118
int x1_dimension = (int)(x->dimensions[0]);
119119
double* restrict x_vals = (double*)(x->vals);
120120

121+
#pragma omp parallel for schedule(static)
122+
for (int32_t py = 0; py < y1_dimension; py++) {
123+
y_vals[py] = 0.0;
124+
}
125+
121126
#pragma omp parallel for schedule(runtime)
122-
for (int32_t i = 0; i < A1_dimension; i++) {
123-
double y_val = 0.0;
124-
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
125-
int32_t j = A2_crd[jA];
126-
y_val += A_vals[jA] * x_vals[j];
127+
for (int32_t i0 = 0; i0 < ((A1_dimension + 31) / 32); i0++) {
128+
for (int32_t i1 = 0; i1 < 32; i1++) {
129+
int32_t i = i0 * 32 + i1;
130+
if (i >= A1_dimension)
131+
continue;
132+
133+
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
134+
int32_t j = A2_crd[jA];
135+
y_vals[i] = y_vals[i] + A_vals[jA] * x_vals[j];
136+
}
127137
}
128-
y_vals[i] = y_val;
129138
}
130139
return 0;
131140
}
@@ -153,14 +162,23 @@ int evaluate(taco_tensor_t *y, taco_tensor_t *A, taco_tensor_t *x) {
153162
int32_t y_capacity = y1_dimension;
154163
y_vals = (double*)malloc(sizeof(double) * y_capacity);
155164

165+
#pragma omp parallel for schedule(static)
166+
for (int32_t py = 0; py < y_capacity; py++) {
167+
y_vals[py] = 0.0;
168+
}
169+
156170
#pragma omp parallel for schedule(runtime)
157-
for (int32_t i = 0; i < A1_dimension; i++) {
158-
double y_val = 0.0;
159-
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
160-
int32_t j = A2_crd[jA];
161-
y_val += A_vals[jA] * x_vals[j];
171+
for (int32_t i0 = 0; i0 < ((A1_dimension + 31) / 32); i0++) {
172+
for (int32_t i1 = 0; i1 < 32; i1++) {
173+
int32_t i = i0 * 32 + i1;
174+
if (i >= A1_dimension)
175+
continue;
176+
177+
for (int32_t jA = A2_pos[i]; jA < A2_pos[(i + 1)]; jA++) {
178+
int32_t j = A2_crd[jA];
179+
y_vals[i] = y_vals[i] + A_vals[jA] * x_vals[j];
180+
}
162181
}
163-
y_vals[i] = y_val;
164182
}
165183

166184
y->vals = (uint8_t*)y_vals;
@@ -218,12 +236,12 @@ int pack_A(taco_tensor_t *A, int* A_COO1_pos, int* A_COO1_crd, int* A_COO2_crd,
218236
jA_COO++;
219237
}
220238
if (A_capacity <= jA) {
221-
A_vals = (double*)realloc(A_vals, sizeof(double) * (A_capacity * 2));
239+
A_vals = (double*)realloc(A_vals, sizeof(double) * A_capacity * 2);
222240
A_capacity *= 2;
223241
}
224242
A_vals[jA] = A_COO_val;
225243
if (A2_crd_size <= jA) {
226-
A2_crd = (int32_t*)realloc(A2_crd, sizeof(int32_t) * (A2_crd_size * 2));
244+
A2_crd = (int32_t*)realloc(A2_crd, sizeof(int32_t) * A2_crd_size * 2);
227245
A2_crd_size *= 2;
228246
}
229247
A2_crd[jA] = j;
@@ -294,12 +312,12 @@ int unpack(int** y_COO1_pos_ptr, int** y_COO1_crd_ptr, double** y_COO_vals_ptr,
294312

295313
for (int32_t i = 0; i < y1_dimension; i++) {
296314
if (y_COO_capacity <= iy_COO) {
297-
y_COO_vals = (double*)realloc(y_COO_vals, sizeof(double) * (y_COO_capacity * 2));
315+
y_COO_vals = (double*)realloc(y_COO_vals, sizeof(double) * y_COO_capacity * 2);
298316
y_COO_capacity *= 2;
299317
}
300318
y_COO_vals[iy_COO] = y_vals[i];
301319
if (y_COO1_crd_size <= iy_COO) {
302-
y_COO1_crd = (int32_t*)realloc(y_COO1_crd, sizeof(int32_t) * (y_COO1_crd_size * 2));
320+
y_COO1_crd = (int32_t*)realloc(y_COO1_crd, sizeof(int32_t) * y_COO1_crd_size * 2);
303321
y_COO1_crd_size *= 2;
304322
}
305323
y_COO1_crd[iy_COO] = i;

‎javascripts/default-schedules.js

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -74,7 +74,7 @@ var default_GPU_schedules = {
7474
},
7575
{
7676
command: "split",
77-
parameters: ["fpos1", "warp", "fpos2", 216]
77+
parameters: ["fpos1", "warp", "fpos2", 256]
7878
},
7979
{
8080
command: "split",

‎javascripts/demo.js

Lines changed: 29 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -5,7 +5,6 @@ function demo() {
55
tensorOrders: {},
66
error: "",
77
indices: [],
8-
accesses: []
98
},
109
schedule: [],
1110
output: {
@@ -31,7 +30,9 @@ function demo() {
3130
}
3231
},
3332
updateScheduleView: function() {
33+
console.log(model.schedule);
3434
model.removeInvalidIndices();
35+
model.removeInvalidAccesses();
3536
model.scheduleView(0);
3637
},
3738
addOutputView: function(newView) {
@@ -65,7 +66,6 @@ function demo() {
6566
try {
6667
model.input.tensorOrders = parser.parse(expression);
6768
model.input.indices = [...new Set(parser_indices.parse(expression))];
68-
model.input.accesses = [...new Set(parser_accesses.parse(expression))];
6969
model.input.error = "";
7070
for (t in model.input.tensorOrders) {
7171
if (model.input.tensorOrders[t] < 0) {
@@ -126,7 +126,7 @@ function demo() {
126126
},
127127
resetSchedule: function() {
128128
model.schedule = [];
129-
model.updateScheduleView();
129+
model.updateScheduleView();
130130
},
131131
addScheduleRow: function() {
132132
model.schedule.push({command: "", parameters: []});
@@ -209,7 +209,24 @@ function demo() {
209209
}
210210
}
211211
},
212+
removeInvalidAccesses: function() {
213+
for (var row = 0; row < model.schedule.length; ++row) {
214+
for (var index = 0; index < model.schedule[row]["parameters"].length; ++index) {
215+
var command = model.schedule[row]["command"];
216+
var value = model.schedule[row]["parameters"][index];
217+
if (model.isParameterType(command, index, "access dropdown")
218+
&& !model.input.tensorOrders.hasOwnProperty(value)) {
219+
model.schedule[row]["parameters"][index] = "";
220+
model.updateInferred(row, command, index, "");
221+
}
222+
}
223+
}
224+
},
212225
isParameterType: function(command, index, parameterType) {
226+
if (command === "reorder") {
227+
return parameterType === "index dropdown";
228+
}
229+
213230
return scheduleCommands[command][index] && scheduleCommands[command][index][0] === parameterType;
214231
},
215232
updateInferred: function(row, command, index, value) {
@@ -352,7 +369,7 @@ function demo() {
352369
var hideTables = function() {
353370
$("#tblFormats").hide();
354371
$("#tblSchedule").hide();
355-
tblScheduleView.clear();
372+
model.resetSchedule()
356373
};
357374
tblFormatsView.timerEvent = setTimeout(hideTables, timeout);
358375
} else {
@@ -698,10 +715,13 @@ function demo() {
698715
// a dropdown where user can choose from argument tensors
699716
function accessDropdown(parameterName, inputId, input) {
700717
var parameter = dropdown(parameterName, inputId, input);
701-
for (var access of model.input.accesses) {
702-
parameter += "<li><a>";
703-
parameter += access;
704-
parameter += "</a></li>";
718+
for (var access in model.input.tensorOrders) {
719+
if (model.input.tensorOrders[access] > 0
720+
&& model.input.expression.indexOf(access) > model.input.expression.indexOf("=")) {
721+
parameter += "<li><a>";
722+
parameter += access;
723+
parameter += "</a></li>";
724+
}
705725
}
706726
parameter += "</ul></div></li>";
707727
return parameter;
@@ -907,7 +927,7 @@ function demo() {
907927

908928
$("#txtExpr").keyup(function() {
909929
model.setInput($("#txtExpr").val());
910-
model.resetSchedule();
930+
model.updateScheduleView();
911931
});
912932

913933
var panelKernelsView = {

0 commit comments

Comments
 (0)