Hi all,
I got a error from a opencl test. The test code is as following:
source code:
__kernel void test_fn(__global uint *src, __global uint4 *results) {
const size_t SIZE = 128;
int tid = get_global_id(0);
if (tid*4 >= SIZE)
return;
__private uint sPrivateStorage[SIZE];
for (size_t i=0 ; i<SIZE ; i++)
sPrivateStorage[i] = src[i];
results[tid] = vload4(tid, sPrivateStorage);
}
and the error message is as following:
"error: variable length arrays are not supported in OpenCL
__private uint sPrivateStorage[SIZE];"
The problem is that 'isArraySizeVLA' function calls
'VerifyIntegerConstantExpression' with 'S.LangOpts.GNUMode'. The OpenCL
Language standard option does not have this option. It causes a error
because 'AllowFold' is not set to 1. In order to fix this error, I think
that we need to insert 'GNUMode' to OpenCL LANGSTANDARD or modify
sligtly checking code. How do you feel about this? I can not guarantee
the side effect of 'GNUMode' with OpenCL. I have attached a simple patch
as reference. If there is something wrong, please let me know.
Thanks,
JinGu Kang
Index: include/clang/Frontend/LangStandards.def
===================================================================
--- include/clang/Frontend/LangStandards.def (revision 195964)
+++ include/clang/Frontend/LangStandards.def (working copy)
@@ -118,13 +118,13 @@
// OpenCL
LANGSTANDARD(opencl, "cl",
"OpenCL 1.0",
- LineComment | C99 | Digraphs | HexFloat)
+ LineComment | C99 | Digraphs | HexFloat | GNUMode)
LANGSTANDARD(opencl11, "CL1.1",
"OpenCL 1.1",
- LineComment | C99 | Digraphs | HexFloat)
+ LineComment | C99 | Digraphs | HexFloat | GNUMode)
LANGSTANDARD(opencl12, "CL1.2",
"OpenCL 1.2",
- LineComment | C99 | Digraphs | HexFloat)
+ LineComment | C99 | Digraphs | HexFloat | GNUMode)
// CUDA
LANGSTANDARD(cuda, "cuda",
Index: lib/Sema/SemaExpr.cpp
===================================================================
--- lib/Sema/SemaExpr.cpp (revision 195964)
+++ lib/Sema/SemaExpr.cpp (working copy)
@@ -10926,7 +10926,7 @@
Notes.clear();
}
- if (!Folded || !AllowFold) {
+ if (!Folded || (!AllowFold && !getLangOpts().OpenCL)) {
if (!Diagnoser.Suppress) {
Diagnoser.diagnoseNotICE(*this, DiagLoc, E->getSourceRange());
for (unsigned I = 0, N = Notes.size(); I != N; ++I)
_______________________________________________
cfe-commits mailing list
[email protected]
http://lists.cs.uiuc.edu/mailman/listinfo/cfe-commits
