mirror of
https://github.com/KhronosGroup/OpenCL-CTS.git
synced 2026-03-26 00:39:03 +00:00
Fix stack-use-after-scope crash in conversions (#1358)
The way that program sources were being constructed involved capturing pointers to strings that were allocated on the stack, and then trying to use them outside of that scope. This change uses a stringstream defined in the outer scope to build the program instead.
This commit is contained in:
@@ -38,6 +38,7 @@
|
|||||||
#include <sys/param.h>
|
#include <sys/param.h>
|
||||||
#endif
|
#endif
|
||||||
|
|
||||||
|
#include <sstream>
|
||||||
#include <stdarg.h>
|
#include <stdarg.h>
|
||||||
#include <stdio.h>
|
#include <stdio.h>
|
||||||
#include <string.h>
|
#include <string.h>
|
||||||
@@ -1559,84 +1560,40 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat,
|
|||||||
cl_program program;
|
cl_program program;
|
||||||
char testName[256];
|
char testName[256];
|
||||||
int error = 0;
|
int error = 0;
|
||||||
const char **strings;
|
|
||||||
size_t stringCount = 0;
|
std::ostringstream source;
|
||||||
|
if (outType == kdouble || inType == kdouble)
|
||||||
|
source << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
||||||
|
|
||||||
// Create the program. This is a bit complicated because we are trying to avoid byte and short stores.
|
// Create the program. This is a bit complicated because we are trying to avoid byte and short stores.
|
||||||
if (0 == vectorSize)
|
if (0 == vectorSize)
|
||||||
{
|
{
|
||||||
|
// Create the type names.
|
||||||
char inName[32];
|
char inName[32];
|
||||||
char outName[32];
|
char outName[32];
|
||||||
const char *programSource[] =
|
|
||||||
{
|
|
||||||
"", // optional pragma
|
|
||||||
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
|
|
||||||
"{\n"
|
|
||||||
" size_t i = get_global_id(0);\n"
|
|
||||||
" dest[i] = src[i];\n"
|
|
||||||
"}\n"
|
|
||||||
};
|
|
||||||
stringCount = sizeof(programSource) / sizeof(programSource[0]);
|
|
||||||
strings = programSource;
|
|
||||||
|
|
||||||
if (outType == kdouble || inType == kdouble)
|
|
||||||
programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
|
||||||
|
|
||||||
//create the type name
|
|
||||||
strncpy(inName, gTypeNames[inType], sizeof(inName));
|
strncpy(inName, gTypeNames[inType], sizeof(inName));
|
||||||
strncpy(outName, gTypeNames[outType], sizeof(outName));
|
strncpy(outName, gTypeNames[outType], sizeof(outName));
|
||||||
sprintf(testName, "test_implicit_%s_%s", outName, inName);
|
sprintf(testName, "test_implicit_%s_%s", outName, inName);
|
||||||
vlog("Building implicit %s -> %s conversion test\n", gTypeNames[inType], gTypeNames[outType]);
|
|
||||||
|
source << "__kernel void " << testName << "( __global " << inName
|
||||||
|
<< " *src, __global " << outName << " *dest )\n";
|
||||||
|
source << "{\n";
|
||||||
|
source << " size_t i = get_global_id(0);\n";
|
||||||
|
source << " dest[i] = src[i];\n";
|
||||||
|
source << "}\n";
|
||||||
|
|
||||||
|
vlog("Building implicit %s -> %s conversion test\n", gTypeNames[inType],
|
||||||
|
gTypeNames[outType]);
|
||||||
fflush(stdout);
|
fflush(stdout);
|
||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
int vectorSizetmp = vectorSizes[vectorSize];
|
int vectorSizetmp = vectorSizes[vectorSize];
|
||||||
|
|
||||||
|
// Create the type names.
|
||||||
char convertString[128];
|
char convertString[128];
|
||||||
char inName[32];
|
char inName[32];
|
||||||
char outName[32];
|
char outName[32];
|
||||||
const char *programSource[] =
|
|
||||||
{
|
|
||||||
"", // optional pragma
|
|
||||||
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
|
|
||||||
"{\n"
|
|
||||||
" size_t i = get_global_id(0);\n"
|
|
||||||
" dest[i] = ", convertString, "( src[i] );\n"
|
|
||||||
"}\n"
|
|
||||||
};
|
|
||||||
const char *programSourceV3[] =
|
|
||||||
{
|
|
||||||
"", // optional pragma
|
|
||||||
"__kernel void ", testName, "( __global ", inName, " *src, __global ", outName, " *dest )\n"
|
|
||||||
"{\n"
|
|
||||||
" size_t i = get_global_id(0);\n"
|
|
||||||
" if( i + 1 < get_global_size(0))\n"
|
|
||||||
" vstore3( ", convertString, "( vload3( i, src)), i, dest );\n"
|
|
||||||
" else\n"
|
|
||||||
" {\n"
|
|
||||||
" ", inName, "3 in;\n"
|
|
||||||
" ", outName, "3 out;\n"
|
|
||||||
" if( 0 == (i & 1) )\n"
|
|
||||||
" in.y = src[3*i+1];\n"
|
|
||||||
" in.x = src[3*i];\n"
|
|
||||||
" out = ", convertString, "( in ); \n"
|
|
||||||
" dest[3*i] = out.x;\n"
|
|
||||||
" if( 0 == (i & 1) )\n"
|
|
||||||
" dest[3*i+1] = out.y;\n"
|
|
||||||
" }\n"
|
|
||||||
"}\n"
|
|
||||||
};
|
|
||||||
stringCount = 3 == vectorSizetmp ? sizeof(programSourceV3) / sizeof(programSourceV3[0]) :
|
|
||||||
sizeof(programSource) / sizeof(programSource[0]);
|
|
||||||
strings = 3 == vectorSizetmp ? programSourceV3 : programSource;
|
|
||||||
|
|
||||||
if (outType == kdouble || inType == kdouble) {
|
|
||||||
programSource[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
|
||||||
programSourceV3[0] = "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
|
|
||||||
}
|
|
||||||
|
|
||||||
//create the type name
|
|
||||||
switch (vectorSizetmp)
|
switch (vectorSizetmp)
|
||||||
{
|
{
|
||||||
case 1:
|
case 1:
|
||||||
@@ -1661,8 +1618,40 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat,
|
|||||||
vlog("Building %s( %s ) test\n", convertString, inName);
|
vlog("Building %s( %s ) test\n", convertString, inName);
|
||||||
break;
|
break;
|
||||||
}
|
}
|
||||||
|
|
||||||
fflush(stdout);
|
fflush(stdout);
|
||||||
|
|
||||||
|
if (vectorSizetmp == 3)
|
||||||
|
{
|
||||||
|
source << "__kernel void " << testName << "( __global " << inName
|
||||||
|
<< " *src, __global " << outName << " *dest )\n";
|
||||||
|
source << "{\n";
|
||||||
|
source << " size_t i = get_global_id(0);\n";
|
||||||
|
source << " if( i + 1 < get_global_size(0))\n";
|
||||||
|
source << " vstore3( " << convertString
|
||||||
|
<< "( vload3( i, src)), i, dest );\n";
|
||||||
|
source << " else\n";
|
||||||
|
source << " {\n";
|
||||||
|
source << " " << inName << "3 in;\n";
|
||||||
|
source << " " << outName << "3 out;\n";
|
||||||
|
source << " if( 0 == (i & 1) )\n";
|
||||||
|
source << " in.y = src[3*i+1];\n";
|
||||||
|
source << " in.x = src[3*i];\n";
|
||||||
|
source << " out = " << convertString << "( in ); \n";
|
||||||
|
source << " dest[3*i] = out.x;\n";
|
||||||
|
source << " if( 0 == (i & 1) )\n";
|
||||||
|
source << " dest[3*i+1] = out.y;\n";
|
||||||
|
source << " }\n";
|
||||||
|
source << "}\n";
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
source << "__kernel void " << testName << "( __global " << inName
|
||||||
|
<< " *src, __global " << outName << " *dest )\n";
|
||||||
|
source << "{\n";
|
||||||
|
source << " size_t i = get_global_id(0);\n";
|
||||||
|
source << " dest[i] = " << convertString << "( src[i] );\n";
|
||||||
|
source << "}\n";
|
||||||
|
}
|
||||||
}
|
}
|
||||||
*outKernel = NULL;
|
*outKernel = NULL;
|
||||||
|
|
||||||
@@ -1671,7 +1660,10 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat,
|
|||||||
flags = "-cl-denorms-are-zero";
|
flags = "-cl-denorms-are-zero";
|
||||||
|
|
||||||
// build it
|
// build it
|
||||||
error = create_single_kernel_helper(gContext, &program, outKernel, (cl_uint)stringCount, strings, testName, flags);
|
std::string sourceString = source.str();
|
||||||
|
const char *programSource = sourceString.c_str();
|
||||||
|
error = create_single_kernel_helper(gContext, &program, outKernel, 1,
|
||||||
|
&programSource, testName, flags);
|
||||||
if (error)
|
if (error)
|
||||||
{
|
{
|
||||||
char buffer[2048] = "";
|
char buffer[2048] = "";
|
||||||
|
|||||||
Reference in New Issue
Block a user