1 /******************************************************************
2 Copyright (c) 2016 The Khronos Group Inc. All Rights Reserved.
3
4 This code is protected by copyright laws and contains material proprietary to the Khronos Group, Inc.
5 This is UNPUBLISHED PROPRIETARY SOURCE CODE that may not be disclosed in whole or in part to
6 third parties, and may not be reproduced, republished, distributed, transmitted, displayed,
7 broadcast or otherwise exploited in any manner without the express prior written permission
8 of Khronos Group. The receipt or possession of this code does not convey any rights to reproduce,
9 disclose, or distribute its contents, or to manufacture, use, or sell anything that it may describe,
10 in whole or in part other than under the terms of the Khronos Adopters Agreement
11 or Khronos Conformance Test Source License Agreement as executed between Khronos and the recipient.
12 ******************************************************************/
13
14 #include "testBase.h"
15 #include "types.hpp"
16
17 #include <sstream>
18 #include <string>
19
20 template<typename T>
test_fmath(cl_device_id deviceID,cl_context context,cl_command_queue queue,const char * spvName,const char * funcName,const char * Tname,bool fast_math,std::vector<T> & h_lhs,std::vector<T> & h_rhs)21 int test_fmath(cl_device_id deviceID,
22 cl_context context,
23 cl_command_queue queue,
24 const char *spvName,
25 const char *funcName,
26 const char *Tname,
27 bool fast_math,
28 std::vector<T> &h_lhs,
29 std::vector<T> &h_rhs)
30 {
31
32 if(std::string(Tname).find("double") != std::string::npos) {
33 if(!is_extension_available(deviceID, "cl_khr_fp64")) {
34 log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n");
35 return 0;
36 }
37 }
38 cl_int err = CL_SUCCESS;
39 int num = (int)h_lhs.size();
40 size_t bytes = num * sizeof(T);
41
42 clMemWrapper lhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err);
43 SPIRV_CHECK_ERROR(err, "Failed to create lhs buffer");
44
45 err = clEnqueueWriteBuffer(queue, lhs, CL_TRUE, 0, bytes, &h_lhs[0], 0, NULL, NULL);
46 SPIRV_CHECK_ERROR(err, "Failed to copy to lhs buffer");
47
48 clMemWrapper rhs = clCreateBuffer(context, CL_MEM_READ_ONLY, bytes, NULL, &err);
49 SPIRV_CHECK_ERROR(err, "Failed to create rhs buffer");
50
51 err = clEnqueueWriteBuffer(queue, rhs, CL_TRUE, 0, bytes, &h_rhs[0], 0, NULL, NULL);
52 SPIRV_CHECK_ERROR(err, "Failed to copy to rhs buffer");
53
54 std::string kernelStr;
55
56 {
57 std::stringstream kernelStream;
58
59 if (is_double<T>::value) {
60 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
61 } else if (sizeof(T) == sizeof(cl_half)) {
62 kernelStream << "#pragma OPENCL EXTENSION cl_khr_fp16 : enable\n";
63 }
64
65 kernelStream << "#define spirv_fadd(a, b) (a) + (b) \n";
66 kernelStream << "#define spirv_fsub(a, b) (a) - (b) \n";
67 kernelStream << "#define spirv_fmul(a, b) (a) * (b) \n";
68 kernelStream << "#define spirv_fdiv(a, b) (a) / (b) \n";
69 kernelStream << "#define spirv_frem(a, b) fmod(a, b) \n";
70 kernelStream << "#define spirv_fmod(a, b) copysign(fmod(a,b),b) \n";
71 kernelStream << "#define T " << Tname << "\n";
72 kernelStream << "#define FUNC spirv_" << funcName << "\n";
73 kernelStream << "__kernel void fmath_cl(__global T *out, \n";
74 kernelStream << "const __global T *lhs, const __global T *rhs) \n";
75 kernelStream << "{ \n";
76 kernelStream << " int id = get_global_id(0); \n";
77 kernelStream << " out[id] = FUNC(lhs[id], rhs[id]); \n";
78 kernelStream << "} \n";
79 kernelStr = kernelStream.str();
80 }
81
82 size_t kernelLen = kernelStr.size();
83 const char *kernelBuf = kernelStr.c_str();
84
85 const char *options = fast_math ? "-cl-fast-relaxed-math" : NULL;
86
87 std::vector<T> h_ref(num);
88
89 {
90 // Run the cl kernel for reference results
91 clProgramWrapper prog;
92 clKernelWrapper kernel;
93 err = create_single_kernel_helper(context, &prog, &kernel, 1,
94 &kernelBuf, "fmath_cl");
95 SPIRV_CHECK_ERROR(err, "Failed to create cl kernel");
96
97 clMemWrapper ref = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
98 SPIRV_CHECK_ERROR(err, "Failed to create ref buffer");
99
100 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &ref);
101 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
102
103 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
104 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
105
106 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
107 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
108
109 size_t global = num;
110 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
111 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
112
113 err = clEnqueueReadBuffer(queue, ref, CL_TRUE, 0, bytes, &h_ref[0], 0, NULL, NULL);
114 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
115 }
116
117 clProgramWrapper prog;
118 err = get_program_with_il(prog, deviceID, context, spvName);
119 SPIRV_CHECK_ERROR(err, "Failed to build program");
120
121 clKernelWrapper kernel = clCreateKernel(prog, "fmath_spv", &err);
122 SPIRV_CHECK_ERROR(err, "Failed to create spv kernel");
123
124 clMemWrapper res = clCreateBuffer(context, CL_MEM_READ_WRITE, bytes, NULL, &err);
125 SPIRV_CHECK_ERROR(err, "Failed to create res buffer");
126
127 err = clSetKernelArg(kernel, 0, sizeof(cl_mem), &res);
128 SPIRV_CHECK_ERROR(err, "Failed to set arg 0");
129
130 err = clSetKernelArg(kernel, 1, sizeof(cl_mem), &lhs);
131 SPIRV_CHECK_ERROR(err, "Failed to set arg 1");
132
133 err = clSetKernelArg(kernel, 2, sizeof(cl_mem), &rhs);
134 SPIRV_CHECK_ERROR(err, "Failed to set arg 2");
135
136 size_t global = num;
137 err = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0, NULL, NULL);
138 SPIRV_CHECK_ERROR(err, "Failed to enqueue cl kernel");
139
140 std::vector<T> h_res(num);
141 err = clEnqueueReadBuffer(queue, res, CL_TRUE, 0, bytes, &h_res[0], 0, NULL, NULL);
142 SPIRV_CHECK_ERROR(err, "Failed to read from ref");
143
144 for (int i = 0; i < num; i++) {
145 if (h_res[i] != h_ref[i]) {
146 log_error("Values do not match at location %d\n", i);
147 return -1;
148 }
149 }
150 return 0;
151 }
152
153 #define TEST_FMATH_FUNC(TYPE, FUNC, MODE) \
154 TEST_SPIRV_FUNC(op_##FUNC##_##TYPE##_##MODE) \
155 { \
156 if (sizeof(cl_##TYPE) == 2) { \
157 PASSIVE_REQUIRE_FP16_SUPPORT(deviceID); \
158 } \
159 const int num = 1 << 20; \
160 std::vector<cl_##TYPE> lhs(num); \
161 std::vector<cl_##TYPE> rhs(num); \
162 \
163 RandomSeed seed(gRandomSeed); \
164 \
165 for (int i = 0; i < num; i++) { \
166 lhs[i] = genrandReal<cl_##TYPE>(seed); \
167 rhs[i] = genrandReal<cl_##TYPE>(seed); \
168 } \
169 \
170 const char *mode = #MODE; \
171 return test_fmath(deviceID, context, queue, \
172 #FUNC "_" #TYPE, \
173 #FUNC, \
174 #TYPE, \
175 mode[0] == 'f', \
176 lhs, rhs); \
177 }
178
179 #define TEST_FMATH_MODE(TYPE, MODE) \
180 TEST_FMATH_FUNC(TYPE, fadd, MODE) \
181 TEST_FMATH_FUNC(TYPE, fsub, MODE) \
182 TEST_FMATH_FUNC(TYPE, fmul, MODE) \
183 TEST_FMATH_FUNC(TYPE, fdiv, MODE) \
184 TEST_FMATH_FUNC(TYPE, frem, MODE) \
185 TEST_FMATH_FUNC(TYPE, fmod, MODE) \
186
187 #define TEST_FMATH_TYPE(TYPE) \
188 TEST_FMATH_MODE(TYPE, regular) \
189 TEST_FMATH_MODE(TYPE, fast) \
190
191 TEST_FMATH_TYPE(float)
192 TEST_FMATH_TYPE(double)
193
194 TEST_FMATH_TYPE(float4)
195 TEST_FMATH_TYPE(double2)
196
197 TEST_FMATH_TYPE(half)
198