ahmed/OpenCL-CTS
1
0
Fork 0
mirror of https://github.com/KhronosGroup/OpenCL-CTS.git synced 2026-03-19 06:09:01 +00:00
Code Issues Packages Projects Releases Wiki Activity

Fix aliasing issue for test_half (#45) (#1004)

Browse Source 
OpenCL-C code must respect aliasing rules as in C99.
Therefore, accessing data from multiple pointer sizes
must all point to character type to keep strict aliasing
rules. If not, when the compiler performs optimization,
the dataflow was broken leading to failure.
We also fix the formatting and refactor few things.
This commit is contained in:
julienhascoet
2020-11-20 15:07:14 +01:00
committed by GitHub
parent 05d5331509
commit f84fcb055b
1 changed files with 27 additions and 6 deletions
Download Patch File Download Diff File Expand all files Collapse all files

33
test_conformance/half/Test_vLoadHalf.cpp
View File

@@ -170,17 +170,38 @@ int Test_vLoadHalf_private( cl_device_id device, bool aligned )
};
const char *source_local2[] = {
"__kernel void test( const __global half *p, __global float", vector_size_name_extensions[vectorSize], " *f )\n"
"#define VECTOR_LEN (",
vector_size_names[vectorSize],
"/",
align_divisors[vectorSize],
")\n"
"#define ALIGN_TYPE ",
align_types[vectorSize],
"\n"
"__kernel void test( const __global half *p, __global float",
vector_size_name_extensions[vectorSize],
" *f )\n"
"{\n"
" __local ", align_types[vectorSize], " data[", local_buf_size, "/", align_divisors[vectorSize], "];\n"
" __local uchar data[",
local_buf_size,
"/",
align_divisors[vectorSize],
"*sizeof(ALIGN_TYPE)] ",
"__attribute__((aligned(sizeof(ALIGN_TYPE))));\n"
" __local half* hdata_p = (__local half*) data;\n"
" __global ", align_types[vectorSize], "* i_p = (__global ", align_types[vectorSize],"*)p;\n"
" __global ALIGN_TYPE* i_p = (__global ALIGN_TYPE*)p;\n"
" size_t i = get_global_id(0);\n"
" size_t lid = get_local_id(0);\n"
" int k;\n"
" for (k=0; k<",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"; k++)\n"
" data[lid*",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"+k] = i_p[i*",vector_size_names[vectorSize],"/",align_divisors[vectorSize],"+k];\n"
" f[i] = vload", aligned ? "a" : "", "_half",vector_size_name_extensions[vectorSize],"( lid, hdata_p );\n"
" for (k=0; k<VECTOR_LEN; k++)\n"
" *(__local ",
"ALIGN_TYPE*)&(data[(lid*VECTOR_LEN+k)*sizeof(ALIGN_TYPE)]) = ",
"i_p[i*VECTOR_LEN+k];\n"
" f[i] = vload",
aligned ? "a" : "",
"_half",
vector_size_name_extensions[vectorSize],
"( lid, hdata_p );\n"
"}\n"
};