diff --git a/test_conformance/conversions/test_conversions.cpp b/test_conformance/conversions/test_conversions.cpp index e8e572e6..d489e28a 100644 --- a/test_conformance/conversions/test_conversions.cpp +++ b/test_conformance/conversions/test_conversions.cpp @@ -38,6 +38,7 @@ #include #endif +#include #include #include #include @@ -1559,84 +1560,40 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat, cl_program program; char testName[256]; 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. if (0 == vectorSize) { + // Create the type names. char inName[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(outName, gTypeNames[outType], sizeof(outName)); 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); } else { int vectorSizetmp = vectorSizes[vectorSize]; + // Create the type names. char convertString[128]; char inName[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) { case 1: @@ -1661,8 +1618,40 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat, vlog("Building %s( %s ) test\n", convertString, inName); break; } - 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; @@ -1671,7 +1660,10 @@ static cl_program MakeProgram( Type outType, Type inType, SaturationMode sat, flags = "-cl-denorms-are-zero"; // 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) { char buffer[2048] = "";