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

Added cl_khr_fp16 extension support for test_vec_type_hint from basic (#1724)

Browse Source 
* Added cl_khr_fp16 extension support for test_vec_type_hint from basic (issue #142, basic)

* Added correction to fix casting problem
This commit is contained in:
Marcin Hajder
2023-06-13 08:39:22 +02:00
committed by GitHub
parent 475a37abbf
commit 095091bc57
1 changed files with 76 additions and 58 deletions
Download Patch File Download Diff File Expand all files Collapse all files

134
test_conformance/basic/test_vec_type_hint.cpp
View File

@@ -13,28 +13,27 @@
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. // limitations under the License.
// //
#include "harness/compat.h"
#include <stdio.h> #include <stdio.h>
#include <stdlib.h> #include <stdlib.h>
#include <string.h> #include <string.h>
#include <sys/types.h> #include <sys/types.h>
#include <sys/stat.h> #include <sys/stat.h>
#include <vector>
#include "procs.h" #include "procs.h"
#include "harness/conversions.h" #include "harness/conversions.h"
#include "harness/typeWrappers.h" #include "harness/typeWrappers.h"
static const char *sample_kernel = { static const char *sample_kernel = {
"%s\n" // optional pragma string "%s\n"
"__kernel __attribute__((vec_type_hint(%s%s))) void sample_test(__global int *src, __global int *dst)\n" "__kernel __attribute__((vec_type_hint(%s%s))) void sample_test(__global "
"{\n" "int *src, __global int *dst)\n"
" int tid = get_global_id(0);\n" "{\n"
" dst[tid] = src[tid];\n" " int tid = get_global_id(0);\n"
"\n" " dst[tid] = src[tid];\n"
"}\n" "\n"
"}\n"
}; };
int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
@@ -42,66 +41,85 @@ int test_vec_type_hint(cl_device_id deviceID, cl_context context, cl_command_que
int error; int error;
int vec_type_index, vec_size_index; int vec_type_index, vec_size_index;
ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt,
const char *size_names[] = {"", "2", "4", "8", "16"}; kLong, kULong, kFloat, kHalf, kDouble };
char *program_source; const char *size_names[] = { "", "2", "4", "8", "16" };
std::vector<char> program_source(4096);
program_source = (char*)malloc(sizeof(char)*4096); for (vec_type_index = 0;
vec_type_index < sizeof(vecType) / sizeof(vecType[0]); vec_type_index++)
{
for (vec_type_index=0; vec_type_index<10; vec_type_index++) { if (vecType[vec_type_index] == kHalf
if (vecType[vec_type_index] == kDouble) { && !is_extension_available(deviceID, "cl_khr_fp16"))
if (!is_extension_available(deviceID, "cl_khr_fp64")) { {
log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); log_info(
continue; "Extension cl_khr_fp16 not supported; skipping half tests.\n");
continue;
}
else if (vecType[vec_type_index] == kDouble
&& !is_extension_available(deviceID, "cl_khr_fp64"))
{
log_info(
"Extension cl_khr_fp64 not supported; skipping double tests.\n");
continue;
}
else if ((vecType[vec_type_index] == kLong
|| vecType[vec_type_index] == kULong)
&& !gHasLong)
{
log_info(
"Extension cl_khr_int64 not supported; skipping long tests.\n");
continue;
} }
log_info("Testing doubles.\n");
}
if (vecType[vec_type_index] == kLong || vecType[vec_type_index] == kULong) for (vec_size_index = 0; vec_size_index < 5; vec_size_index++)
{ {
if (!gHasLong) clProgramWrapper program;
{ clKernelWrapper kernel;
log_info("Extension cl_khr_int64 not supported; skipping long tests.\n"); clMemWrapper in, out;
continue; size_t global[] = { 1, 1, 1 };
}
}
for (vec_size_index=0; vec_size_index<5; vec_size_index++) { log_info("Testing __attribute__((vec_type_hint(%s%s))...\n",
clProgramWrapper program; get_explicit_type_name(vecType[vec_type_index]),
clKernelWrapper kernel; size_names[vec_size_index]);
clMemWrapper in, out; char extension[128] = { 0 };
size_t global[] = {1,1,1}; if (vecType[vec_type_index] == kDouble)
std::snprintf(extension, sizeof(extension),
"#pragma OPENCL EXTENSION cl_khr_fp64 : enable");
else if (vecType[vec_type_index] == kHalf)
std::snprintf(extension, sizeof(extension),
"#pragma OPENCL EXTENSION cl_khr_fp16 : enable");
log_info("Testing __attribute__((vec_type_hint(%s%s))...\n", get_explicit_type_name(vecType[vec_type_index]), size_names[vec_size_index]); sprintf(program_source.data(), sample_kernel, extension,
get_explicit_type_name(vecType[vec_type_index]),
size_names[vec_size_index]);
program_source[0] = '\0'; const char *src = &program_source.front();
sprintf(program_source, sample_kernel, error = create_single_kernel_helper(context, &program, &kernel, 1,
(vecType[vec_type_index] == kDouble) ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", &src, "sample_test");
get_explicit_type_name(vecType[vec_type_index]), size_names[vec_size_index]); test_error(error, "create_single_kernel_helper failed");
error = create_single_kernel_helper( context, &program, &kernel, 1, (const char**)&program_source, "sample_test" ); in = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_int) * 10,
if( error != 0 ) NULL, &error);
return error; test_error(error, "clCreateBuffer failed");
out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_int) * 10,
NULL, &error);
test_error(error, "clCreateBuffer failed");
in = clCreateBuffer(context, CL_MEM_READ_ONLY, sizeof(cl_int)*10, NULL, &error); error = clSetKernelArg(kernel, 0, sizeof(in), &in);
test_error(error, "clCreateBuffer failed"); test_error(error, "clSetKernelArg failed");
out = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_int)*10, NULL, &error); error = clSetKernelArg(kernel, 1, sizeof(out), &out);
test_error(error, "clCreateBuffer failed"); test_error(error, "clSetKernelArg failed");
error = clSetKernelArg(kernel, 0, sizeof(in), &in); error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, NULL,
test_error(error, "clSetKernelArg failed"); 0, NULL, NULL);
error = clSetKernelArg(kernel, 1, sizeof(out), &out); test_error(error, "clEnqueueNDRangeKernel failed");
test_error(error, "clSetKernelArg failed");
error = clEnqueueNDRangeKernel(queue, kernel, 3, NULL, global, NULL, 0, NULL, NULL); error = clFinish(queue);
test_error(error, "clEnqueueNDRangeKernel failed"); test_error(error, "clFinish failed");
}
error = clFinish(queue);
test_error(error, "clFinish failed");
}
} }
free(program_source);
return 0; return 0;
} }