Merge pull request #65 from pavanky/kernel_fixes

Timmy · Timmy · commit fc692eb46dfe · 2015-02-16T13:14:25.000-06:00
Always cast get_*_* to uints before casting to vectors.
diff --git a/src/library/blas/gens/gemm.c b/src/library/blas/gens/gemm.c
@@ -560,7 +560,7 @@ blockGen(
         kgenAddStmt(ctx, tmp);
     }
     else {
-        sprintf(globalIdB, "get_global_id(%d)", 1-i);
+        sprintf(globalIdB, "(uint)get_global_id(%d)", 1-i);
     }
 
     if (!(isColMajA || isColMajB)) {
@@ -758,7 +758,7 @@ subgGen(
     vecLenA = gset.tileA.vecLen;
 
     // channel offset based coordinate
-    ksprintf(&exprK, "( get_group_id(0)*%lu + k )", staggered/vecLenA*vecLenA);
+    ksprintf(&exprK, "( (uint)(get_group_id(0))*%lu + k )", staggered/vecLenA*vecLenA);
 
     // starting code generation--------------------------------------------------
     pCtx = createKgenContext(pBuf, buflen, true);

Original file line number	Diff line number	Diff line change
`@@ -560,7 +560,7 @@ blockGen(`
`560`	`560`	`kgenAddStmt(ctx, tmp);`
`561`	`561`	`}`
`562`	`562`	`else {`
`563`		`- sprintf(globalIdB, "get_global_id(%d)", 1-i);`
	`563`	`+ sprintf(globalIdB, "(uint)get_global_id(%d)", 1-i);`
`564`	`564`	`}`
`565`	`565`
`566`	`566`	`if (!(isColMajA \|\| isColMajB)) {`
`@@ -758,7 +758,7 @@ subgGen(`
`758`	`758`	`vecLenA = gset.tileA.vecLen;`
`759`	`759`
`760`	`760`	`// channel offset based coordinate`
`761`		`- ksprintf(&exprK, "( get_group_id(0)%lu + k )", staggered/vecLenAvecLenA);`
	`761`	`+ ksprintf(&exprK, "( (uint)(get_group_id(0))%lu + k )", staggered/vecLenAvecLenA);`
`762`	`762`
`763`	`763`	`// starting code generation--------------------------------------------------`
`764`	`764`	`pCtx = createKgenContext(pBuf, buflen, true);`