1 /******************************************************************
2 Copyright (c) 2020 The Khronos Group Inc. All Rights Reserved.
3
4 This code is protected by copyright laws and contains material proprietary to
5 the Khronos Group, Inc. This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not
6 be disclosed in whole or in part to third parties, and may not be reproduced,
7 republished, distributed, transmitted, displayed, broadcast or otherwise
8 exploited in any manner without the express prior written permission of Khronos
9 Group. The receipt or possession of this code does not convey any rights to
10 reproduce, disclose, or distribute its contents, or to manufacture, use, or sell
11 anything that it may describe, in whole or in part other than under the terms of
12 the Khronos Adopters Agreement or Khronos Conformance Test Source License
13 Agreement as executed between Khronos and the recipient.
14 ******************************************************************/
15
16 #include "testBase.h"
17 #include "types.hpp"
18
19
20 template <typename T>
run_case(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * name,T init_buffer,T spec_constant_value,T final_value,bool use_spec_constant)21 int run_case(cl_device_id deviceID, cl_context context, cl_command_queue queue,
22 const char *name, T init_buffer, T spec_constant_value,
23 T final_value, bool use_spec_constant)
24 {
25 clProgramWrapper prog;
26 cl_int err = CL_SUCCESS;
27 if (use_spec_constant)
28 {
29 spec_const new_spec_const =
30 spec_const(101, sizeof(T), &spec_constant_value);
31
32 err =
33 get_program_with_il(prog, deviceID, context, name, new_spec_const);
34 }
35 else
36 {
37 err = get_program_with_il(prog, deviceID, context, name);
38 }
39 SPIRV_CHECK_ERROR(err, "Failed to build program");
40
41 clKernelWrapper kernel = clCreateKernel(prog, "spec_const_kernel", &err);
42 SPIRV_CHECK_ERROR(err, "Failed to create kernel");
43 size_t bytes = sizeof(T);
44 clMemWrapper output_buffer =
45 clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_USE_HOST_PTR, bytes,
46 &init_buffer, &err);
47 SPIRV_CHECK_ERROR(err, "Failed to create output_buffer");
48
49 err = clSetKernelArg(kernel, 0, sizeof(clMemWrapper), &output_buffer);
50 SPIRV_CHECK_ERROR(err, "Failed to set kernel argument output_buffer");
51
52 size_t work_size = 1;
53 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &work_size, NULL, 0,
54 NULL, NULL);
55 SPIRV_CHECK_ERROR(err, "Failed to enqueue kernel");
56
57 T device_results = 0;
58 err = clEnqueueReadBuffer(queue, output_buffer, CL_TRUE, 0, bytes,
59 &device_results, 0, NULL, NULL);
60 SPIRV_CHECK_ERROR(err, "Failed to copy from output_buffer");
61 T reference = 0;
62 use_spec_constant ? reference = final_value : reference = init_buffer;
63 if (device_results != reference)
64 {
65 log_error("Values do not match. Expected %d obtained %d\n", reference,
66 device_results);
67 err = -1;
68 }
69 return err;
70 }
71
72 template <typename T>
test_spec_constant(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * name,T init_buffer,T spec_constant_value,T final_value)73 int test_spec_constant(cl_device_id deviceID, cl_context context,
74 cl_command_queue queue, const char *name, T init_buffer,
75 T spec_constant_value, T final_value)
76 {
77 if (std::string(name).find("double") != std::string::npos)
78 {
79 if (!is_extension_available(deviceID, "cl_khr_fp64"))
80 {
81 log_info("Extension cl_khr_fp64 not supported; skipping double "
82 "tests.\n");
83 return TEST_SKIPPED_ITSELF;
84 }
85 }
86 if (std::string(name).find("half") != std::string::npos)
87 {
88 if (!is_extension_available(deviceID, "cl_khr_fp16"))
89 {
90 log_info("Extension cl_khr_fp16 not supported; skipping half "
91 "tests.\n");
92 return TEST_SKIPPED_ITSELF;
93 }
94 }
95 cl_int err = CL_SUCCESS;
96 err = run_case<T>(deviceID, context, queue, name, init_buffer,
97 spec_constant_value, final_value, false);
98 err |= run_case<T>(deviceID, context, queue, name, init_buffer,
99 spec_constant_value, final_value, true);
100
101 if (err == CL_SUCCESS)
102 {
103 return TEST_PASS;
104 }
105 else
106 {
107 return TEST_FAIL;
108 }
109 }
110
111
112 #define TEST_SPEC_CONSTANT(NAME, type, init_buffer, spec_constant_value) \
113 TEST_SPIRV_FUNC_VERSION(op_spec_constant_##NAME##_simple, Version(2, 2)) \
114 { \
115 type init_value = init_buffer; \
116 type final_value = init_value + spec_constant_value; \
117 return test_spec_constant( \
118 deviceID, context, queue, "op_spec_constant_" #NAME "_simple", \
119 init_value, (type)spec_constant_value, final_value); \
120 }
121
122 // type name, type, value init, spec constant value
123 TEST_SPEC_CONSTANT(uint, cl_uint, 25, 43)
124 TEST_SPEC_CONSTANT(uchar, cl_uchar, 19, 4)
125 TEST_SPEC_CONSTANT(ushort, cl_ushort, 6000, 3000)
126 TEST_SPEC_CONSTANT(ulong, cl_ulong, 9223372036854775000UL, 200)
127 TEST_SPEC_CONSTANT(float, cl_float, 1.5, -3.7)
128 TEST_SPEC_CONSTANT(half, cl_half, 1, 2)
129 TEST_SPEC_CONSTANT(double, cl_double, 14534.53453, 1.53453)
130
131 // Boolean tests
132 // documenation: 'If a specialization constant is a boolean
133 // constant, spec_value should be a pointer to a cl_uchar value'
134
135 TEST_SPIRV_FUNC_VERSION(op_spec_constant_true_simple, Version(2, 2))
136 {
137 // 1-st ndrange init_value is expected value (no change)
138 // 2-nd ndrange sets spec const to 'false' so value = value + 1
139 cl_uchar value = (cl_uchar)7;
140 cl_uchar init_value = value;
141 cl_uchar final_value = value + 1;
142 return test_spec_constant<cl_uchar>(deviceID, context, queue,
143 "op_spec_constant_true_simple",
144 init_value, 0, final_value);
145 }
146
147 TEST_SPIRV_FUNC_VERSION(op_spec_constant_false_simple, Version(2, 2))
148 {
149 // 1-st ndrange init_value is expected value (no change)
150 // 2-nd ndrange sets spec const to 'true' so value = value + 1
151 cl_uchar value = (cl_uchar)7;
152 cl_uchar init_value = value;
153 cl_uchar final_value = value + 1;
154 return test_spec_constant<cl_uchar>(deviceID, context, queue,
155 "op_spec_constant_false_simple",
156 init_value, 1, final_value);
157 }
158