/*
* Copyright © Microsoft Corporation
*
* Permission is hereby granted, free of charge, to any person obtaining a
* copy of this software and associated documentation files (the "Software"),
* to deal in the Software without restriction, including without limitation
* the rights to use, copy, modify, merge, publish, distribute, sublicense,
* and/or sell copies of the Software, and to permit persons to whom the
* Software is furnished to do so, subject to the following conditions:
*
* The above copyright notice and this permission notice (including the next
* paragraph) shall be included in all copies or substantial portions of the
* Software.
*
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
* IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
* FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
* THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
* LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
* FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
* IN THE SOFTWARE.
*/
#include <stdio.h>
#include <stdint.h>
#include <stdexcept>
#include <vector>
#include <directx/d3d12.h>
#include <dxgi1_4.h>
#include <gtest/gtest.h>
#include <wrl.h>
#include "compute_test.h"
using std::vector;
TEST_F(ComputeTest, runtime_memcpy)
{
struct shift { uint8_t val; uint8_t shift; uint16_t ret; };
const char *kernel_source =
"struct shift { uchar val; uchar shift; ushort ret; };\n\
__kernel void main_test(__global struct shift *inout)\n\
{\n\
uint id = get_global_id(0);\n\
uint id2 = id + get_global_id(1);\n\
struct shift lc[4] = { { 0, 0, 0 }, { 0, 0, 0 }, { 0, 0, 0 }, { 0, 0, 0 }};\n\
lc[id] = inout[id];\n\
inout[id2].ret = (ushort) lc[id2].val << (ushort) lc[id2].shift;\n\
}\n";
auto inout = ShaderArg<struct shift>({
{ 0x10, 1, 0xffff },
{ 0x20, 2, 0xffff },
{ 0x30, 3, 0xffff },
{ 0x40, 4, 0xffff },
},
SHADER_ARG_INOUT);
const uint16_t expected[] = { 0x20, 0x80, 0x180, 0x400 };
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i].ret, expected[i]);
}
TEST_F(ComputeTest, two_global_arrays)
{
const char *kernel_source =
"__kernel void main_test(__global uint *g1, __global uint *g2)\n\
{\n\
uint idx = get_global_id(0);\n\
g1[idx] -= g2[idx];\n\
}\n";
auto g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto g2 = ShaderArg<uint32_t>({ 1, 2, 3, 4 }, SHADER_ARG_INPUT);
const uint32_t expected[] = {
9, 18, 27, 36
};
run_shader(kernel_source, g1.size(), 1, 1, g1, g2);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected[i]);
}
/* Disabled until saturated conversions from f32->i64 fixed (mesa/mesa#3824) */
TEST_F(ComputeTest, DISABLED_i64tof32)
{
const char *kernel_source =
"__kernel void main_test(__global long *out, __constant long *in)\n\
{\n\
__local float tmp[12];\n\
uint idx = get_global_id(0);\n\
tmp[idx] = in[idx];\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
out[idx] = tmp[idx + get_global_id(1)];\n\
}\n";
auto in = ShaderArg<int64_t>({ 0x100000000LL,
-0x100000000LL,
0x7fffffffffffffffLL,
0x4000004000000000LL,
0x4000003fffffffffLL,
0x4000004000000001LL,
-1,
-0x4000004000000000LL,
-0x4000003fffffffffLL,
-0x4000004000000001LL,
0,
INT64_MIN },
SHADER_ARG_INPUT);
auto out = ShaderArg<int64_t>(std::vector<int64_t>(12, 0xdeadbeed), SHADER_ARG_OUTPUT);
const int64_t expected[] = {
0x100000000LL,
-0x100000000LL,
0x7fffffffffffffffLL,
0x4000000000000000LL,
0x4000000000000000LL,
0x4000008000000000LL,
-1,
-0x4000000000000000LL,
-0x4000000000000000LL,
-0x4000008000000000LL,
0,
INT64_MIN,
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ((int64_t)out[i], expected[i]);
}
}
TEST_F(ComputeTest, two_constant_arrays)
{
const char *kernel_source =
"__kernel void main_test(__constant uint *c1, __global uint *g1, __constant uint *c2)\n\
{\n\
uint idx = get_global_id(0);\n\
g1[idx] -= c1[idx] + c2[idx];\n\
}\n";
auto g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto c1 = ShaderArg<uint32_t>({ 1, 2, 3, 4 }, SHADER_ARG_INPUT);
auto c2 = ShaderArg<uint32_t>(std::vector<uint32_t>(16384, 5), SHADER_ARG_INPUT);
const uint32_t expected[] = {
4, 13, 22, 31
};
run_shader(kernel_source, g1.size(), 1, 1, c1, g1, c2);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected[i]);
}
TEST_F(ComputeTest, null_constant_ptr)
{
const char *kernel_source =
"__kernel void main_test(__global uint *g1, __constant uint *c1)\n\
{\n\
__constant uint fallback[] = {2, 3, 4, 5};\n\
__constant uint *c = c1 ? c1 : fallback;\n\
uint idx = get_global_id(0);\n\
g1[idx] -= c[idx];\n\
}\n";
auto g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto c1 = ShaderArg<uint32_t>({ 1, 2, 3, 4 }, SHADER_ARG_INPUT);
const uint32_t expected1[] = {
9, 18, 27, 36
};
run_shader(kernel_source, g1.size(), 1, 1, g1, c1);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected1[i]);
const uint32_t expected2[] = {
8, 17, 26, 35
};
g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto c2 = NullShaderArg();
run_shader(kernel_source, g1.size(), 1, 1, g1, c2);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected2[i]);
}
/* This test seems to fail on older versions of WARP. */
TEST_F(ComputeTest, DISABLED_null_global_ptr)
{
const char *kernel_source =
"__kernel void main_test(__global uint *g1, __global uint *g2)\n\
{\n\
__constant uint fallback[] = {2, 3, 4, 5};\n\
uint idx = get_global_id(0);\n\
g1[idx] -= g2 ? g2[idx] : fallback[idx];\n\
}\n";
auto g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto g2 = ShaderArg<uint32_t>({ 1, 2, 3, 4 }, SHADER_ARG_INPUT);
const uint32_t expected1[] = {
9, 18, 27, 36
};
run_shader(kernel_source, g1.size(), 1, 1, g1, g2);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected1[i]);
const uint32_t expected2[] = {
8, 17, 26, 35
};
g1 = ShaderArg<uint32_t>({ 10, 20, 30, 40 }, SHADER_ARG_INOUT);
auto g2null = NullShaderArg();
run_shader(kernel_source, g1.size(), 1, 1, g1, g2null);
for (int i = 0; i < g1.size(); ++i)
EXPECT_EQ(g1[i], expected2[i]);
}
TEST_F(ComputeTest, ret_constant_ptr)
{
struct s { uint64_t ptr; uint32_t val; };
const char *kernel_source =
"struct s { __constant uint *ptr; uint val; };\n\
__kernel void main_test(__global struct s *out, __constant uint *in)\n\
{\n\
__constant uint foo[] = { 1, 2 };\n\
uint idx = get_global_id(0);\n\
if (idx == 0)\n\
out[idx].ptr = foo;\n\
else\n\
out[idx].ptr = in;\n\
out[idx].val = out[idx].ptr[idx];\n\
}\n";
auto out = ShaderArg<struct s>(std::vector<struct s>(2, {0xdeadbeefdeadbeef, 0}), SHADER_ARG_OUTPUT);
auto in = ShaderArg<uint32_t>({ 3, 4 }, SHADER_ARG_INPUT);
const uint32_t expected_val[] = {
1, 4
};
const uint64_t expected_ptr[] = {
2ull << 32, 1ull << 32
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].val, expected_val[i]);
EXPECT_EQ(out[i].ptr, expected_ptr[i]);
}
}
TEST_F(ComputeTest, ret_global_ptr)
{
struct s { uint64_t ptr; uint32_t val; };
const char *kernel_source =
"struct s { __global uint *ptr; uint val; };\n\
__kernel void main_test(__global struct s *out, __global uint *in1, __global uint *in2)\n\
{\n\
uint idx = get_global_id(0);\n\
out[idx].ptr = idx ? in2 : in1;\n\
out[idx].val = out[idx].ptr[idx];\n\
}\n";
auto out = ShaderArg<struct s>(std::vector<struct s>(2, {0xdeadbeefdeadbeef, 0}), SHADER_ARG_OUTPUT);
auto in1 = ShaderArg<uint32_t>({ 1, 2 }, SHADER_ARG_INPUT);
auto in2 = ShaderArg<uint32_t>({ 3, 4 }, SHADER_ARG_INPUT);
const uint32_t expected_val[] = {
1, 4
};
const uint64_t expected_ptr[] = {
1ull << 32, 2ull << 32
};
run_shader(kernel_source, out.size(), 1, 1, out, in1, in2);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].val, expected_val[i]);
EXPECT_EQ(out[i].ptr, expected_ptr[i]);
}
}
TEST_F(ComputeTest, ret_local_ptr)
{
struct s { uint64_t ptr; };
const char *kernel_source =
"struct s { __local uint *ptr; };\n\
__kernel void main_test(__global struct s *out)\n\
{\n\
__local uint tmp[2];\n\
uint idx = get_global_id(0);\n\
tmp[idx] = idx;\n\
out[idx].ptr = &tmp[idx];\n\
}\n";
auto out = ShaderArg<struct s>(std::vector<struct s>(2, { 0xdeadbeefdeadbeef }), SHADER_ARG_OUTPUT);
const uint64_t expected_ptr[] = {
0, 4,
};
run_shader(kernel_source, out.size(), 1, 1, out);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].ptr, expected_ptr[i]);
}
}
TEST_F(ComputeTest, ret_private_ptr)
{
struct s { uint64_t ptr; uint32_t value; };
const char *kernel_source =
"struct s { __private uint *ptr; uint value; };\n\
__kernel void main_test(__global struct s *out)\n\
{\n\
uint tmp[2] = {1, 2};\n\
uint idx = get_global_id(0);\n\
out[idx].ptr = &tmp[idx];\n\
out[idx].value = *out[idx].ptr;\n\
}\n";
auto out = ShaderArg<struct s>(std::vector<struct s>(2, { 0xdeadbeefdeadbeef }), SHADER_ARG_OUTPUT);
const uint64_t expected_ptr[] = {
0, 4,
};
const uint32_t expected_value[] = {
1, 2
};
run_shader(kernel_source, out.size(), 1, 1, out);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].ptr, expected_ptr[i]);
}
}
TEST_F(ComputeTest, globals_8bit)
{
const char *kernel_source =
"__kernel void main_test(__global unsigned char *inout)\n\
{\n\
uint idx = get_global_id(0);\n\
inout[idx] = inout[idx] + 1;\n\
}\n";
auto inout = ShaderArg<uint8_t> ({ 100, 110, 120, 130 }, SHADER_ARG_INOUT);
const uint8_t expected[] = {
101, 111, 121, 131
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, globals_16bit)
{
const char *kernel_source =
"__kernel void main_test(__global unsigned short *inout)\n\
{\n\
uint idx = get_global_id(0);\n\
inout[idx] = inout[idx] + 1;\n\
}\n";
auto inout = ShaderArg<uint16_t> ({ 10000, 10010, 10020, 10030 }, SHADER_ARG_INOUT);
const uint16_t expected[] = {
10001, 10011, 10021, 10031
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, DISABLED_globals_64bit)
{
/* Test disabled, because we need a fixed version of WARP that hasn't
been officially shipped yet */
const char *kernel_source =
"__kernel void main_test(__global unsigned long *inout)\n\
{\n\
uint idx = get_global_id(0);\n\
inout[idx] = inout[idx] + 1;\n\
}\n";
uint64_t base = 1ull << 50;
auto inout = ShaderArg<uint64_t>({ base, base + 10, base + 20, base + 30 },
SHADER_ARG_INOUT);
const uint64_t expected[] = {
base + 1, base + 11, base + 21, base + 31
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, built_ins_global_id)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = get_global_id(0);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
0, 1, 2, 3
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, built_ins_global_id_rmw)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
uint id = get_global_id(0);\n\
output[id] = output[id] * (id + 1);\n\
}\n";
auto inout = ShaderArg<uint32_t>({0x00000001, 0x10000001, 0x00020002, 0x04010203},
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0x00000001, 0x20000002, 0x00060006, 0x1004080c
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, types_float_basics)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = (uint)((float)get_global_id(0) + 1.5f);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 2, 3, 4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, DISABLED_types_double_basics)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = (uint)((double)get_global_id(0) + 1.5);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 2, 3, 4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, types_short_basics)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = (uint)((short)get_global_id(0) + (short)1);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 2, 3, 4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, types_char_basics)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = (uint)((char)get_global_id(0) + (char)1);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 2, 3, 4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, types_if_statement)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
int idx = get_global_id(0);\n\
if (idx > 0)\n\
output[idx] = ~idx;\n\
else\n\
output[0] = 0xff;\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
0xff, ~1u, ~2u, ~3u
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, types_do_while_loop)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
int value = 1;\n\
int i = 1, n = get_global_id(0);\n\
do {\n\
value *= i++;\n\
} while (i <= n);\n\
output[n] = value;\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(5, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 1, 1*2, 1*2*3, 1*2*3*4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, types_for_loop)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
int value = 1;\n\
int n = get_global_id(0);\n\
for (int i = 1; i <= n; ++i)\n\
value *= i;\n\
output[n] = value;\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(5, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
1, 1, 1*2, 1*2*3, 1*2*3*4
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, DISABLED_complex_types_local_array_long)
{
const char *kernel_source =
"__kernel void main_test(__global ulong *inout)\n\
{\n\
ushort tmp[] = {\n\
get_global_id(1) + 0x00000000,\n\
get_global_id(1) + 0x10000001,\n\
get_global_id(1) + 0x20000020,\n\
get_global_id(1) + 0x30000300,\n\
};\n\
uint idx = get_global_id(0);\n\
inout[idx] = tmp[idx];\n\
}\n";
auto inout = ShaderArg<uint64_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
const uint16_t expected[] = {
0x00000000, 0x10000001, 0x20000020, 0x30000300,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, complex_types_local_array_short)
{
const char *kernel_source =
"__kernel void main_test(__global ushort *inout)\n\
{\n\
ushort tmp[] = {\n\
get_global_id(1) + 0x00,\n\
get_global_id(1) + 0x10,\n\
get_global_id(1) + 0x20,\n\
get_global_id(1) + 0x30,\n\
};\n\
uint idx = get_global_id(0);\n\
inout[idx] = tmp[idx];\n\
}\n";
auto inout = ShaderArg<uint16_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
const uint16_t expected[] = {
0x00, 0x10, 0x20, 0x30,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, complex_types_local_array_struct_vec_float_misaligned)
{
const char *kernel_source =
"struct has_vecs { uchar c; ushort s; float2 f; };\n\
__kernel void main_test(__global uint *inout)\n\
{\n\
struct has_vecs tmp[] = {\n\
{ 10 + get_global_id(0), get_global_id(1), { 10.0f, 1.0f } },\n\
{ 19 + get_global_id(0), get_global_id(1), { 20.0f, 4.0f } },\n\
{ 28 + get_global_id(0), get_global_id(1), { 30.0f, 9.0f } },\n\
{ 37 + get_global_id(0), get_global_id(1), { 40.0f, 16.0f } },\n\
};\n\
uint idx = get_global_id(0);\n\
uint mul = (tmp[idx].c + tmp[idx].s) * trunc(tmp[idx].f[0]);\n\
inout[idx] = mul + trunc(tmp[idx].f[1]);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
const uint16_t expected[] = { 101, 404, 909, 1616 };
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, complex_types_local_array)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
uint tmp[] = {\n\
get_global_id(1) + 0x00,\n\
get_global_id(1) + 0x10,\n\
get_global_id(1) + 0x20,\n\
get_global_id(1) + 0x30,\n\
};\n\
uint idx = get_global_id(0);\n\
inout[idx] = tmp[idx];\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
const uint32_t expected[] = {
0x00, 0x10, 0x20, 0x30,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, complex_types_global_struct_array)
{
struct two_vals { uint32_t add; uint32_t mul; };
const char *kernel_source =
"struct two_vals { uint add; uint mul; };\n\
__kernel void main_test(__global struct two_vals *in_out)\n\
{\n\
uint id = get_global_id(0);\n\
in_out[id].add = in_out[id].add + id;\n\
in_out[id].mul = in_out[id].mul * id;\n\
}\n";
auto inout = ShaderArg<struct two_vals>({ { 8, 8 }, { 16, 16 }, { 64, 64 }, { 65536, 65536 } },
SHADER_ARG_INOUT);
const struct two_vals expected[] = {
{ 8 + 0, 8 * 0 },
{ 16 + 1, 16 * 1 },
{ 64 + 2, 64 * 2 },
{ 65536 + 3, 65536 * 3 }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].add, expected[i].add);
EXPECT_EQ(inout[i].mul, expected[i].mul);
}
}
TEST_F(ComputeTest, complex_types_global_uint2)
{
struct uint2 { uint32_t x; uint32_t y; };
const char *kernel_source =
"__kernel void main_test(__global uint2 *inout)\n\
{\n\
uint id = get_global_id(0);\n\
inout[id].x = inout[id].x + id;\n\
inout[id].y = inout[id].y * id;\n\
}\n";
auto inout = ShaderArg<struct uint2>({ { 8, 8 }, { 16, 16 }, { 64, 64 }, { 65536, 65536 } },
SHADER_ARG_INOUT);
const struct uint2 expected[] = {
{ 8 + 0, 8 * 0 },
{ 16 + 1, 16 * 1 },
{ 64 + 2, 64 * 2 },
{ 65536 + 3, 65536 * 3 }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].x, expected[i].x);
EXPECT_EQ(inout[i].y, expected[i].y);
}
}
TEST_F(ComputeTest, complex_types_global_ushort2)
{
struct ushort2 { uint16_t x; uint16_t y; };
const char *kernel_source =
"__kernel void main_test(__global ushort2 *inout)\n\
{\n\
uint id = get_global_id(0);\n\
inout[id].x = inout[id].x + id;\n\
inout[id].y = inout[id].y * id;\n\
}\n";
auto inout = ShaderArg<struct ushort2>({ { 8, 8 }, { 16, 16 }, { 64, 64 },
{ (uint16_t)65536, (uint16_t)65536 } },
SHADER_ARG_INOUT);
const struct ushort2 expected[] = {
{ 8 + 0, 8 * 0 },
{ 16 + 1, 16 * 1 },
{ 64 + 2, 64 * 2 },
{ (uint16_t)(65536 + 3), (uint16_t)(65536 * 3) }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].x, expected[i].x);
EXPECT_EQ(inout[i].y, expected[i].y);
}
}
TEST_F(ComputeTest, complex_types_global_uchar3)
{
struct uchar3 { uint8_t x; uint8_t y; uint8_t z; uint8_t pad; };
const char *kernel_source =
"__kernel void main_test(__global uchar3 *inout)\n\
{\n\
uint id = get_global_id(0);\n\
inout[id].x = inout[id].x + id;\n\
inout[id].y = inout[id].y * id;\n\
inout[id].z = inout[id].y + inout[id].x;\n\
}\n";
auto inout = ShaderArg<struct uchar3>({ { 8, 8, 8 }, { 16, 16, 16 }, { 64, 64, 64 }, { 255, 255, 255 } },
SHADER_ARG_INOUT);
const struct uchar3 expected[] = {
{ 8 + 0, 8 * 0, (8 + 0) + (8 * 0) },
{ 16 + 1, 16 * 1, (16 + 1) + (16 * 1) },
{ 64 + 2, 64 * 2, (64 + 2) + (64 * 2) },
{ (uint8_t)(255 + 3), (uint8_t)(255 * 3), (uint8_t)((255 + 3) + (255 * 3)) }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].x, expected[i].x);
EXPECT_EQ(inout[i].y, expected[i].y);
EXPECT_EQ(inout[i].z, expected[i].z);
}
}
TEST_F(ComputeTest, complex_types_constant_uchar3)
{
struct uchar3 { uint8_t x; uint8_t y; uint8_t z; uint8_t pad; };
const char *kernel_source =
"__kernel void main_test(__global uchar3 *out, __constant uchar3 *in)\n\
{\n\
uint id = get_global_id(0);\n\
out[id].x = in[id].x + id;\n\
out[id].y = in[id].y * id;\n\
out[id].z = out[id].y + out[id].x;\n\
}\n";
auto in = ShaderArg<struct uchar3>({ { 8, 8, 8 }, { 16, 16, 16 }, { 64, 64, 64 }, { 255, 255, 255 } },
SHADER_ARG_INPUT);
auto out = ShaderArg<struct uchar3>(std::vector<struct uchar3>(4, { 0xff, 0xff, 0xff }),
SHADER_ARG_OUTPUT);
const struct uchar3 expected[] = {
{ 8 + 0, 8 * 0, (8 + 0) + (8 * 0) },
{ 16 + 1, 16 * 1, (16 + 1) + (16 * 1) },
{ 64 + 2, 64 * 2, (64 + 2) + (64 * 2) },
{ (uint8_t)(255 + 3), (uint8_t)(255 * 3), (uint8_t)((255 + 3) + (255 * 3)) }
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].x, expected[i].x);
EXPECT_EQ(out[i].y, expected[i].y);
EXPECT_EQ(out[i].z, expected[i].z);
}
}
TEST_F(ComputeTest, complex_types_global_uint8)
{
struct uint8 {
uint32_t s0; uint32_t s1; uint32_t s2; uint32_t s3;
uint32_t s4; uint32_t s5; uint32_t s6; uint32_t s7;
};
const char *kernel_source =
"__kernel void main_test(__global uint8 *inout)\n\
{\n\
uint id = get_global_id(0);\n\
inout[id].s01234567 = inout[id].s01234567 * 2;\n\
}\n";
auto inout = ShaderArg<struct uint8>({ { 1, 2, 3, 4, 5, 6, 7, 8 } },
SHADER_ARG_INOUT);
const struct uint8 expected[] = {
{ 2, 4, 6, 8, 10, 12, 14, 16 }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].s0, expected[i].s0);
EXPECT_EQ(inout[i].s1, expected[i].s1);
EXPECT_EQ(inout[i].s2, expected[i].s2);
EXPECT_EQ(inout[i].s3, expected[i].s3);
EXPECT_EQ(inout[i].s4, expected[i].s4);
EXPECT_EQ(inout[i].s5, expected[i].s5);
EXPECT_EQ(inout[i].s6, expected[i].s6);
EXPECT_EQ(inout[i].s7, expected[i].s7);
}
}
TEST_F(ComputeTest, complex_types_local_ulong16)
{
struct ulong16 {
uint64_t values[16];
};
const char *kernel_source =
R"(__kernel void main_test(__global ulong16 *inout)
{
__local ulong16 local_array[2];
uint id = get_global_id(0);
local_array[id] = inout[id];
barrier(CLK_LOCAL_MEM_FENCE);
inout[id] = local_array[0] * 2;
})";
auto inout = ShaderArg<struct ulong16>({ { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } },
SHADER_ARG_INOUT);
const struct ulong16 expected[] = {
{ 2, 4, 6, 8, 10, 12, 14, 16, 18, 20, 22, 24, 26, 28, 30 }
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
for (int j = 0; j < 16; ++j) {
EXPECT_EQ(inout[i].values[j], expected[i].values[j]);
}
}
}
TEST_F(ComputeTest, complex_types_constant_uint8)
{
struct uint8 {
uint32_t s0; uint32_t s1; uint32_t s2; uint32_t s3;
uint32_t s4; uint32_t s5; uint32_t s6; uint32_t s7;
};
const char *kernel_source =
"__kernel void main_test(__global uint8 *out, __constant uint8 *in)\n\
{\n\
uint id = get_global_id(0);\n\
out[id].s01234567 = in[id].s01234567 * 2;\n\
}\n";
auto in = ShaderArg<struct uint8>({ { 1, 2, 3, 4, 5, 6, 7, 8 } },
SHADER_ARG_INPUT);
auto out = ShaderArg<struct uint8>({ { 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff, 0xff } },
SHADER_ARG_INOUT);
const struct uint8 expected[] = {
{ 2, 4, 6, 8, 10, 12, 14, 16 }
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].s0, expected[i].s0);
EXPECT_EQ(out[i].s1, expected[i].s1);
EXPECT_EQ(out[i].s2, expected[i].s2);
EXPECT_EQ(out[i].s3, expected[i].s3);
EXPECT_EQ(out[i].s4, expected[i].s4);
EXPECT_EQ(out[i].s5, expected[i].s5);
EXPECT_EQ(out[i].s6, expected[i].s6);
EXPECT_EQ(out[i].s7, expected[i].s7);
}
}
TEST_F(ComputeTest, DISABLED_complex_types_const_array)
{
/* DISABLED because current release versions of WARP either return
* rubbish from reads or crash: they are not prepared to handle
* non-float global constants */
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
const uint foo[] = { 100, 101, 102, 103 };\n\
output[get_global_id(0)] = foo[get_global_id(0) % 4];\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
100, 101, 102, 103
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, mem_access_load_store_ordering)
{
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
uint foo[4];\n\
foo[0] = 0x11111111;\n\
foo[1] = 0x22222222;\n\
foo[2] = 0x44444444;\n\
foo[3] = 0x88888888;\n\
foo[get_global_id(1)] -= 0x11111111; // foo[0] = 0 \n\
foo[0] += get_global_id(0); // foo[0] = tid\n\
foo[foo[get_global_id(1)]] = get_global_id(0); // foo[tid] = tid\n\
output[get_global_id(0)] = foo[get_global_id(0)]; // output[tid] = tid\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint16_t expected[] = {
0, 1, 2, 3
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, DISABLED_two_const_arrays)
{
/* DISABLED because current release versions of WARP either return
* rubbish from reads or crash: they are not prepared to handle
* non-float global constants */
const char *kernel_source =
"__kernel void main_test(__global uint *output)\n\
{\n\
uint id = get_global_id(0);\n\
uint foo[4] = {100, 101, 102, 103};\n\
uint bar[4] = {1, 2, 3, 4};\n\
output[id] = foo[id] * bar[id];\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
100, 202, 306, 412
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, imod_pos)
{
const char *kernel_source =
"__kernel void main_test(__global int *inout)\n\
{\n\
inout[get_global_id(0)] = inout[get_global_id(0)] % 3;\n\
}\n";
auto inout = ShaderArg<int32_t>({ -4, -3, -2, -1, 0, 1, 2, 3, 4 },
SHADER_ARG_INOUT);
const int32_t expected[] = {
-1, 0, -2, -1, 0, 1, 2, 0, 1
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, imod_neg)
{
const char *kernel_source =
"__kernel void main_test(__global int *inout)\n\
{\n\
inout[get_global_id(0)] = inout[get_global_id(0)] % -3;\n\
}\n";
auto inout = ShaderArg<int32_t>({ -4, -3, -2, -1, 0, 1, 2, 3, 4 },
SHADER_ARG_INOUT);
const int32_t expected[] = {
-1, 0, -2, -1, 0, 1, 2, 0, 1
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, umod)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = inout[get_global_id(0)] % 0xfffffffc;\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0xfffffffa, 0xfffffffb, 0xfffffffc, 0xfffffffd, 0xfffffffe },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0xfffffffa, 0xfffffffb, 0, 1, 2
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, rotate)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = rotate(inout[get_global_id(0)], (uint)get_global_id(0) * 4);\n\
}\n";
auto inout = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0xdeadbeef, 0xeadbeefd, 0xadbeefde, 0xdbeefdea
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, popcount)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = popcount(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 0x1, 0x3, 0x101, 0x110011, ~0u },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0, 1, 2, 2, 4, 32
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, hadd)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = hadd(inout[get_global_id(0)], 1u << 31);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, 0xfffffffc, 0xfffffffd, 0xfffffffe, 0xffffffff },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
(1u << 31) >> 1,
((1u << 31) + 1) >> 1,
((1u << 31) + 2) >> 1,
((1u << 31) + 3) >> 1,
((1ull << 31) + 0xfffffffc) >> 1,
((1ull << 31) + 0xfffffffd) >> 1,
((1ull << 31) + 0xfffffffe) >> 1,
((1ull << 31) + 0xffffffff) >> 1,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, rhadd)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = rhadd(inout[get_global_id(0)], 1u << 31);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, 0xfffffffc, 0xfffffffd, 0xfffffffe, 0xffffffff },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
((1u << 31) + 1) >> 1,
((1u << 31) + 2) >> 1,
((1u << 31) + 3) >> 1,
((1u << 31) + 4) >> 1,
((1ull << 31) + 0xfffffffd) >> 1,
((1ull << 31) + 0xfffffffe) >> 1,
((1ull << 31) + 0xffffffff) >> 1,
((1ull << 31) + (1ull << 32)) >> 1,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, add_sat)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = add_sat(inout[get_global_id(0)], 2u);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0xffffffff - 3, 0xffffffff - 2, 0xffffffff - 1, 0xffffffff },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0xffffffff - 1, 0xffffffff, 0xffffffff, 0xffffffff
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, sub_sat)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = sub_sat(inout[get_global_id(0)], 2u);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3 }, SHADER_ARG_INOUT);
const uint32_t expected[] = {
0, 0, 0, 1
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, mul_hi)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = mul_hi(inout[get_global_id(0)], 1u << 31);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, (1u << 31) }, SHADER_ARG_INOUT);
const uint32_t expected[] = {
0, 0, 1, 1, (1u << 30)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, ldexp_x)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = ldexp(inout[get_global_id(0)], 5);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 2.0f }, SHADER_ARG_INOUT);
const float expected[] = {
ldexp(0.0f, 5), ldexp(0.5f, 5), ldexp(1.0f, 5), ldexp(2.0f, 5)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, ldexp_y)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = ldexp(inout[get_global_id(0)], get_global_id(0));\n\
}\n";
auto inout = ShaderArg<float>({ 0.25f, 0.5f, 0.75f, 1.0f }, SHADER_ARG_INOUT);
const float expected[] = {
ldexp(0.25f, 0), ldexp(0.5f, 1), ldexp(0.75f, 2), ldexp(1.0f, 3)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, frexp_ret)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
int exp;\n\
inout[get_global_id(0)] = frexp(inout[get_global_id(0)], &exp);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
0.0f, 0.5f, 0.5f, 0.75f
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, frexp_exp)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
int exp;\n\
frexp(inout[get_global_id(0)], &exp);\n\
inout[get_global_id(0)] = (float)exp;\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
0.0f, 0.0f, 1.0f, 2.0f
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, clz)
{
const char *kernel_source =
"__kernel void main_test(__global uint *inout)\n\
{\n\
inout[get_global_id(0)] = clz(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 0, 1, 0xffff, (1u << 30), (1u << 31) }, SHADER_ARG_INOUT);
const uint32_t expected[] = {
32, 31, 16, 1, 0
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, sin)
{
struct sin_vals { float in; float clc; float native; };
const char *kernel_source =
"struct sin_vals { float in; float clc; float native; };\n\
__kernel void main_test(__global struct sin_vals *inout)\n\
{\n\
inout[get_global_id(0)].clc = sin(inout[get_global_id(0)].in);\n\
inout[get_global_id(0)].native = native_sin(inout[get_global_id(0)].in);\n\
}\n";
const vector<sin_vals> input = {
{ 0.0f, 0.0f, 0.0f },
{ 1.0f, 0.0f, 0.0f },
{ 2.0f, 0.0f, 0.0f },
{ 3.0f, 0.0f, 0.0f },
};
auto inout = ShaderArg<sin_vals>(input, SHADER_ARG_INOUT);
const struct sin_vals expected[] = {
{ 0.0f, 0.0f, 0.0f },
{ 1.0f, sin(1.0f), sin(1.0f) },
{ 2.0f, sin(2.0f), sin(2.0f) },
{ 3.0f, sin(3.0f), sin(3.0f) },
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_FLOAT_EQ(inout[i].in, inout[i].in);
EXPECT_FLOAT_EQ(inout[i].clc, inout[i].clc);
EXPECT_NEAR(inout[i].clc, inout[i].native, 0.008f); // range from DXIL spec
}
}
TEST_F(ComputeTest, DISABLED_cosh)
{
/* Disabled because of WARP failures, where we fetch incorrect results when
* sourcing from non-float ICBs */
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = cosh(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
cosh(0.0f), cosh(1.0f), cosh(2.0f), cosh(3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, exp)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_exp(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
exp(0.0f), exp(1.0f), exp(2.0f), exp(3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, exp10)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_exp10(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
pow(10.0f, 0.0f), pow(10.0f, 1.0f), pow(10.0f, 2.0f), pow(10.0f, 3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, exp2)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_exp2(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
pow(2.0f, 0.0f), pow(2.0f, 1.0f), pow(2.0f, 2.0f), pow(2.0f, 3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, log)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_log(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
log(0.0f), log(1.0f), log(2.0f), log(3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, log10)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_log10(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
log10(0.0f), log10(1.0f), log10(2.0f), log10(3.0f)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, log2)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = native_log2(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
const float expected[] = {
log(0.0f) / log(2), log(1.0f) / log(2), log(2.0f) / log(2), log(3.0f) / log(2)
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, rint)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = rint(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0.5f, 1.5f, -0.5f, -1.5f, 1.4f }, SHADER_ARG_INOUT);
const float expected[] = {
0.0f, 2.0f, 0.0f, -2.0f, 1.0f,
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, round)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = round(inout[get_global_id(0)]);\n\
}\n";
auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
SHADER_ARG_INOUT);
const float expected[] = {
0.0f, 0.0f, -0.0f, 1.0f, -1.0f, 1.0f, -1.0f
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, arg_by_val)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout, float mul)\n\
{\n\
inout[get_global_id(0)] = inout[get_global_id(0)] * mul;\n\
}\n";
auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
SHADER_ARG_INOUT);
auto mul = ShaderArg<float>(10.0f, SHADER_ARG_INPUT);
const float expected[] = {
0.0f, 3.0f, -3.0f, 5.0f, -5.0f, 11.0f, -11.0f
};
run_shader(kernel_source, inout.size(), 1, 1, inout, mul);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, uint8_by_val)
{
struct uint8 {
uint32_t s0; uint32_t s1; uint32_t s2; uint32_t s3;
uint32_t s4; uint32_t s5; uint32_t s6; uint32_t s7;
};
const char *kernel_source =
"__kernel void main_test(__global uint *out, uint8 val)\n\
{\n\
out[get_global_id(0)] = val.s0 + val.s1 + val.s2 + val.s3 +\n\
val.s4 + val.s5 + val.s6 + val.s7;\n\
}\n";
auto out = ShaderArg<uint32_t>({ 0 }, SHADER_ARG_OUTPUT);
auto val = ShaderArg<struct uint8>({ {0, 1, 2, 3, 4, 5, 6, 7 }}, SHADER_ARG_INPUT);
const uint32_t expected[] = { 0 + 1 + 2 + 3 + 4 + 5 + 6 + 7 };
run_shader(kernel_source, out.size(), 1, 1, out, val);
for (int i = 0; i < out.size(); ++i)
EXPECT_EQ(out[i], expected[i]);
}
TEST_F(ComputeTest, link)
{
const char *foo_src =
"float foo(float in)\n\
{\n\
return in * in;\n\
}\n";
const char *kernel_source =
"float foo(float in);\n\
__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = foo(inout[get_global_id(0)]);\n\
}\n";
std::vector<const char *> srcs = { foo_src, kernel_source };
auto inout = ShaderArg<float>({ 2.0f }, SHADER_ARG_INOUT);
const float expected[] = {
4.0f,
};
run_shader(srcs, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, link_library)
{
const char *bar_src =
"float bar(float in)\n\
{\n\
return in * 5;\n\
}\n";
const char *foo_src =
"float bar(float in);\n\
float foo(float in)\n\
{\n\
return in * bar(in);\n\
}\n";
const char *kernel_source =
"float foo(float in);\n\
__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] = foo(inout[get_global_id(0)]);\n\
}\n";
std::vector<Shader> libraries = {
compile({ bar_src, kernel_source }, {}, true),
compile({ foo_src }, {}, true)
};
Shader exe = link(libraries);
auto inout = ShaderArg<float>({ 2.0f }, SHADER_ARG_INOUT);
const float expected[] = {
20.0f,
};
run_shader(exe, { (unsigned)inout.size(), 1, 1 }, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, localvar)
{
const char *kernel_source =
"__kernel __attribute__((reqd_work_group_size(2, 1, 1)))\n\
void main_test(__global float *inout)\n\
{\n\
__local float2 tmp[2];\n\
tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
}\n";
auto inout = ShaderArg<float>({ 2.0f, 4.0f }, SHADER_ARG_INOUT);
const float expected[] = {
9.0f, 5.0f
};
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, localvar_uchar2)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(2, 1, 1)))\n\
__kernel void main_test(__global uchar *inout)\n\
{\n\
__local uchar2 tmp[2];\n\
tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
}\n";
auto inout = ShaderArg<uint8_t>({ 2, 4 }, SHADER_ARG_INOUT);
const uint8_t expected[] = { 9, 5 };
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, work_group_size_hint)
{
const char *kernel_source =
"__attribute__((work_group_size_hint(2, 1, 1)))\n\
__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = get_local_id(0);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
0, 1, 2, 3
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, reqd_work_group_size)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(2, 1, 1)))\n\
__kernel void main_test(__global uint *output)\n\
{\n\
output[get_global_id(0)] = get_local_id(0);\n\
}\n";
auto output = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
SHADER_ARG_OUTPUT);
const uint32_t expected[] = {
0, 1, 0, 1
};
run_shader(kernel_source, output.size(), 1, 1, output);
for (int i = 0; i < output.size(); ++i)
EXPECT_EQ(output[i], expected[i]);
}
TEST_F(ComputeTest, image)
{
const char* kernel_source =
"__kernel void main_test(read_only image2d_t input, write_only image2d_t output)\n\
{\n\
int2 coords = (int2)(get_global_id(0), get_global_id(1));\n\
write_imagef(output, coords, read_imagef(input, coords));\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, image_two_reads)
{
const char* kernel_source =
"__kernel void main_test(image2d_t image, int is_float, __global float* output)\n\
{\n\
if (is_float)\n\
output[get_global_id(0)] = read_imagef(image, (int2)(0, 0)).x;\n\
else \n\
output[get_global_id(0)] = (float)read_imagei(image, (int2)(0, 0)).x;\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, image_read_write)
{
const char *kernel_source =
R"(__kernel void main_test(read_write image2d_t image)
{
int2 coords = (int2)(get_global_id(0), get_global_id(1));
write_imagef(image, coords, read_imagef(image, coords) + (float4)(1.0f, 1.0f, 1.0f, 1.0f));
})";
Shader shader = compile(std::vector<const char*>({ kernel_source }), { "-cl-std=cl3.0" });
validate(shader);
}
TEST_F(ComputeTest, sampler)
{
const char* kernel_source =
"__kernel void main_test(image2d_t image, sampler_t sampler, __global float* output)\n\
{\n\
output[get_global_id(0)] = read_imagef(image, sampler, (int2)(0, 0)).x;\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, image_dims)
{
const char* kernel_source =
"__kernel void main_test(image2d_t roimage, write_only image2d_t woimage, __global uint* output)\n\
{\n\
output[get_global_id(0)] = get_image_width(roimage);\n\
output[get_global_id(0) + 1] = get_image_width(woimage);\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, image_format)
{
const char* kernel_source =
"__kernel void main_test(image2d_t roimage, write_only image2d_t woimage, __global uint* output)\n\
{\n\
output[get_global_id(0)] = get_image_channel_data_type(roimage);\n\
output[get_global_id(0) + 1] = get_image_channel_order(woimage);\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, image1d_buffer_t)
{
const char* kernel_source =
"__kernel void main_test(read_only image1d_buffer_t input, write_only image1d_buffer_t output)\n\
{\n\
write_imageui(output, get_global_id(0), read_imageui(input, get_global_id(0)));\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, local_ptr)
{
struct uint2 { uint32_t x, y; };
const char *kernel_source =
"__kernel void main_test(__global uint *inout, __local uint2 *tmp)\n\
{\n\
tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 2, 4 }, SHADER_ARG_INOUT);
auto tmp = ShaderArg<struct uint2>(std::vector<struct uint2>(4096), SHADER_ARG_INPUT);
const uint8_t expected[] = { 9, 5 };
run_shader(kernel_source, inout.size(), 1, 1, inout, tmp);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, two_local_ptrs)
{
struct uint2 { uint32_t x, y; };
const char *kernel_source =
"__kernel void main_test(__global uint *inout, __local uint2 *tmp, __local uint *tmp2)\n\
{\n\
tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
tmp2[get_local_id(0)] = get_global_id(0);\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y + tmp2[get_local_id(0) % 2];\n\
}\n";
auto inout = ShaderArg<uint32_t>({ 2, 4 }, SHADER_ARG_INOUT);
auto tmp = ShaderArg<struct uint2>(std::vector<struct uint2>(1024), SHADER_ARG_INPUT);
auto tmp2 = ShaderArg<uint32_t>(std::vector<uint32_t>(1024), SHADER_ARG_INPUT);
const uint8_t expected[] = { 9, 6 };
run_shader(kernel_source, inout.size(), 1, 1, inout, tmp, tmp2);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, int8_to_float)
{
const char *kernel_source =
"__kernel void main_test(__global char* in, __global float* out)\n\
{\n\
uint pos = get_global_id(0);\n\
out[pos] = in[pos] / 100.0f;\n\
}";
auto in = ShaderArg<char>({ 10, 20, 30, 40 }, SHADER_ARG_INPUT);
auto out = ShaderArg<float>(std::vector<float>(4, std::numeric_limits<float>::infinity()), SHADER_ARG_OUTPUT);
const float expected[] = { 0.1f, 0.2f, 0.3f, 0.4f };
run_shader(kernel_source, in.size(), 1, 1, in, out);
for (int i = 0; i < in.size(); ++i)
EXPECT_FLOAT_EQ(out[i], expected[i]);
}
TEST_F(ComputeTest, vec_hint_float4)
{
const char *kernel_source =
"__kernel __attribute__((vec_type_hint(float4))) void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 4);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_FLOAT);
}
TEST_F(ComputeTest, vec_hint_uchar2)
{
const char *kernel_source =
"__kernel __attribute__((vec_type_hint(uchar2))) void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 2);
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_type, CLC_VEC_HINT_TYPE_CHAR);
}
TEST_F(ComputeTest, vec_hint_none)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout)\n\
{\n\
inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
}";
Shader shader = compile({ kernel_source });
EXPECT_EQ(shader.metadata->kernels[0].vec_hint_size, 0);
}
TEST_F(ComputeTest, DISABLED_debug_layer_failure)
{
const char *kernel_source =
"__kernel void main_test(__global float *inout, float mul)\n\
{\n\
inout[get_global_id(0)] = inout[get_global_id(0)] * mul;\n\
}\n";
auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
SHADER_ARG_INOUT);
auto mul = ShaderArg<float>(10.0f, SHADER_ARG_INPUT);
const float expected[] = {
0.0f, 3.0f, -3.0f, 5.0f, -5.0f, 11.0f, -11.0f
};
ComPtr<ID3D12InfoQueue> info_queue;
dev->QueryInterface(info_queue.ReleaseAndGetAddressOf());
if (!info_queue) {
GTEST_SKIP() << "No info queue";
return;
}
info_queue->AddApplicationMessage(D3D12_MESSAGE_SEVERITY_ERROR, "This should cause the test to fail");
run_shader(kernel_source, inout.size(), 1, 1, inout, mul);
for (int i = 0; i < inout.size(); ++i)
EXPECT_FLOAT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, compiler_defines)
{
const char *kernel_source =
"__kernel void main_test(__global int* out)\n\
{\n\
out[0] = OUT_VAL0;\n\
out[1] = __OPENCL_C_VERSION__;\n\
}";
auto out = ShaderArg<int>(std::vector<int>(2, 0), SHADER_ARG_OUTPUT);
CompileArgs compile_args = { 1, 1, 1 };
compile_args.compiler_command_line = { "-DOUT_VAL0=5", "-cl-std=cl" };
std::vector<RawShaderArg *> raw_args = { &out };
run_shader({ kernel_source }, compile_args, out);
EXPECT_EQ(out[0], 5);
EXPECT_EQ(out[1], 100);
}
/* There's a bug in WARP turning atomic_add(ptr, x) into
* atomic_add(ptr, x * 4). Works fine on intel HW.
*/
TEST_F(ComputeTest, DISABLED_global_atomic_add)
{
const char *kernel_source =
"__kernel void main_test(__global int *inout, __global int *old)\n\
{\n\
old[get_global_id(0)] = atomic_add(inout + get_global_id(0), 3);\n\
}\n";
auto inout = ShaderArg<int32_t>({ 2, 4 }, SHADER_ARG_INOUT);
auto old = ShaderArg<int32_t>(std::vector<int32_t>(2, 0xdeadbeef), SHADER_ARG_OUTPUT);
const int32_t expected_inout[] = { 5, 7 };
const int32_t expected_old[] = { 2, 4 };
run_shader(kernel_source, inout.size(), 1, 1, inout, old);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i], expected_inout[i]);
EXPECT_EQ(old[i], expected_old[i]);
}
}
TEST_F(ComputeTest, global_atomic_imin)
{
const char *kernel_source =
"__kernel void main_test(__global int *inout, __global int *old)\n\
{\n\
old[get_global_id(0)] = atomic_min(inout + get_global_id(0), 1);\n\
}\n";
auto inout = ShaderArg<int32_t>({ 0, 2, -1 }, SHADER_ARG_INOUT);
auto old = ShaderArg<int32_t>(std::vector<int32_t>(3, 0xdeadbeef), SHADER_ARG_OUTPUT);
const int32_t expected_inout[] = { 0, 1, -1 };
const int32_t expected_old[] = { 0, 2, -1 };
run_shader(kernel_source, inout.size(), 1, 1, inout, old);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i], expected_inout[i]);
EXPECT_EQ(old[i], expected_old[i]);
}
}
TEST_F(ComputeTest, global_atomic_and_or)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(3, 1, 1)))\n\
__kernel void main_test(__global int *inout)\n\
{\n\
atomic_and(inout, ~(1 << get_global_id(0)));\n\
atomic_or(inout, (1 << (get_global_id(0) + 4)));\n\
}\n";
auto inout = ShaderArg<int32_t>(0xf, SHADER_ARG_INOUT);
const int32_t expected[] = { 0x78 };
run_shader(kernel_source, 3, 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, global_atomic_cmpxchg)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(2, 1, 1)))\n\
__kernel void main_test(__global int *inout)\n\
{\n\
while (atomic_cmpxchg(inout, get_global_id(0), get_global_id(0) + 1) != get_global_id(0))\n\
;\n\
}\n";
auto inout = ShaderArg<int32_t>(0, SHADER_ARG_INOUT);
const int32_t expected_inout[] = { 2 };
run_shader(kernel_source, 2, 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected_inout[i]);
}
TEST_F(ComputeTest, local_atomic_and_or)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(2, 1, 1)))\n\
__kernel void main_test(__global ushort *inout)\n\
{\n\
__local ushort tmp;\n\
atomic_and(&tmp, ~(0xff << (get_global_id(0) * 8)));\n\
atomic_or(&tmp, inout[get_global_id(0)] << (get_global_id(0) * 8));\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
inout[get_global_id(0)] = tmp;\n\
}\n";
auto inout = ShaderArg<uint16_t>({ 2, 4 }, SHADER_ARG_INOUT);
const uint16_t expected[] = { 0x402, 0x402 };
run_shader(kernel_source, inout.size(), 1, 1, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}
TEST_F(ComputeTest, local_atomic_cmpxchg)
{
const char *kernel_source =
"__attribute__((reqd_work_group_size(2, 1, 1)))\n\
__kernel void main_test(__global int *out)\n\
{\n\
__local uint tmp;\n\
tmp = 0;\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
while (atomic_cmpxchg(&tmp, get_global_id(0), get_global_id(0) + 1) != get_global_id(0))\n\
;\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
out[0] = tmp;\n\
}\n";
auto out = ShaderArg<uint32_t>(0xdeadbeef, SHADER_ARG_OUTPUT);
const uint16_t expected[] = { 2 };
run_shader(kernel_source, 2, 1, 1, out);
for (int i = 0; i < out.size(); ++i)
EXPECT_EQ(out[i], expected[i]);
}
TEST_F(ComputeTest, constant_sampler)
{
const char* kernel_source =
"__constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_LINEAR;\n\
__kernel void main_test(read_only image2d_t input, write_only image2d_t output)\n\
{\n\
int2 coordsi = (int2)(get_global_id(0), get_global_id(1));\n\
float2 coordsf = (float2)((float)coordsi.x / get_image_width(input), (float)coordsi.y / get_image_height(input));\n\
write_imagef(output, coordsi, \n\
read_imagef(input, sampler, coordsf) + \n\
read_imagef(input, sampler, coordsf + (float2)(0.1, 0.1)));\n\
}\n";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
EXPECT_EQ(shader.dxil->metadata.num_const_samplers, 1);
}
TEST_F(ComputeTest, hi)
{
const char *kernel_source = R"(
__kernel void main_test(__global char3 *srcA, __global char2 *dst)
{
int tid = get_global_id(0);
char2 tmp = srcA[tid].hi;
dst[tid] = tmp;
})";
Shader shader = compile(std::vector<const char*>({ kernel_source }));
validate(shader);
}
TEST_F(ComputeTest, system_values)
{
const char *kernel_source =
"__kernel void main_test(__global uint* outputs)\n\
{\n\
outputs[0] = get_work_dim();\n\
outputs[1] = get_global_size(0);\n\
outputs[2] = get_local_size(0);\n\
outputs[3] = get_num_groups(0);\n\
outputs[4] = get_group_id(0);\n\
outputs[5] = get_global_offset(0);\n\
outputs[6] = get_global_id(0);\n\
}\n";
auto out = ShaderArg<uint32_t>(std::vector<uint32_t>(6, 0xdeadbeef), SHADER_ARG_OUTPUT);
const uint16_t expected[] = { 3, 1, 1, 1, 0, 0, 0, };
CompileArgs args = { 1, 1, 1 };
Shader shader = compile({ kernel_source });
run_shader(shader, args, out);
for (int i = 0; i < out.size(); ++i)
EXPECT_EQ(out[i], expected[i]);
args.work_props.work_dim = 2;
args.work_props.global_offset_x = 100;
args.work_props.group_id_offset_x = 2;
args.work_props.group_count_total_x = 5;
const uint32_t expected_withoffsets[] = { 2, 5, 1, 5, 2, 100, 102 };
run_shader(shader, args, out);
for (int i = 0; i < out.size(); ++i)
EXPECT_EQ(out[i], expected_withoffsets[i]);
}
TEST_F(ComputeTest, convert_round_sat)
{
const char *kernel_source =
"__kernel void main_test(__global float *f, __global uchar *u)\n\
{\n\
uint idx = get_global_id(0);\n\
u[idx] = convert_uchar_sat_rtp(f[idx]);\n\
}\n";
auto f = ShaderArg<float>({ -1.0f, 1.1f, 20.0f, 255.5f }, SHADER_ARG_INPUT);
auto u = ShaderArg<uint8_t>({ 255, 0, 0, 0 }, SHADER_ARG_OUTPUT);
const uint8_t expected[] = {
0, 2, 20, 255
};
run_shader(kernel_source, f.size(), 1, 1, f, u);
for (int i = 0; i < u.size(); ++i)
EXPECT_EQ(u[i], expected[i]);
}
TEST_F(ComputeTest, convert_round_sat_vec)
{
const char *kernel_source =
"__kernel void main_test(__global float16 *f, __global uchar16 *u)\n\
{\n\
uint idx = get_global_id(0);\n\
u[idx] = convert_uchar16_sat_rtp(f[idx]);\n\
}\n";
auto f = ShaderArg<float>({
-1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f,
-0.5f, 1.9f, 20.0f, 254.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f,
0.0f, 1.3f, 20.0f, 255.1f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f,
-0.0f, 1.5555f, 20.0f, 254.9f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f, -1.0f, 1.1f, 20.0f, 255.5f,
}, SHADER_ARG_INPUT);
auto u = ShaderArg<uint8_t>({
255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0,
255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0,
255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0,
255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0, 255, 0, 0, 0,
}, SHADER_ARG_OUTPUT);
const uint8_t expected[] = {
0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255,
0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255,
0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255,
0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255, 0, 2, 20, 255,
};
run_shader(kernel_source, 4, 1, 1, f, u);
for (int i = 0; i < u.size(); ++i)
EXPECT_EQ(u[i], expected[i]);
}
TEST_F(ComputeTest, convert_char2_uchar2)
{
const char *kernel_source =
"__kernel void main_test( __global char2 *src, __global uchar2 *dest )\n\
{\n\
size_t i = get_global_id(0);\n\
dest[i] = convert_uchar2_sat( src[i] );\n\
}\n";
auto c = ShaderArg<int8_t>({ -127, -4, 0, 4, 126, 127, 16, 32 }, SHADER_ARG_INPUT);
auto u = ShaderArg<uint8_t>({ 99, 99, 99, 99, 99, 99, 99, 99 }, SHADER_ARG_OUTPUT);
const uint8_t expected[] = { 0, 0, 0, 4, 126, 127, 16, 32 };
run_shader(kernel_source, 4, 1, 1, c, u);
for (int i = 0; i < u.size(); i++)
EXPECT_EQ(u[i], expected[i]);
}
TEST_F(ComputeTest, async_copy)
{
const char *kernel_source = R"(
__kernel void main_test( const __global char *src, __global char *dst, __local char *localBuffer, int copiesPerWorkgroup, int copiesPerWorkItem )
{
int i;
for(i=0; i<copiesPerWorkItem; i++)
localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ] = (char)(char)0;
barrier( CLK_LOCAL_MEM_FENCE );
event_t event;
event = async_work_group_copy( (__local char*)localBuffer, (__global const char*)(src+copiesPerWorkgroup*get_group_id(0)), (size_t)copiesPerWorkgroup, 0 );
wait_group_events( 1, &event );
for(i=0; i<copiesPerWorkItem; i++)
dst[ get_global_id( 0 )*copiesPerWorkItem+i ] = localBuffer[ get_local_id( 0 )*copiesPerWorkItem+i ];
})";
Shader shader = compile({ kernel_source });
validate(shader);
}
TEST_F(ComputeTest, packed_struct_global)
{
#pragma pack(push, 1)
struct s { uint8_t uc; uint64_t ul; uint16_t us; };
#pragma pack(pop)
const char *kernel_source =
"struct __attribute__((packed)) s {uchar uc; ulong ul; ushort us; };\n\
__kernel void main_test(__global struct s *inout, global uint *size)\n\
{\n\
uint idx = get_global_id(0);\n\
inout[idx].uc = idx + 1;\n\
inout[idx].ul = ((ulong)(idx + 1 + 0xfbfcfdfe) << 32) | 0x12345678;\n\
inout[idx].us = ((ulong)(idx + 1 + 0xa0) << 8) | 0x12;\n\
*size = sizeof(struct s);\n\
}\n";
auto inout = ShaderArg<struct s>({0, 0, 0}, SHADER_ARG_OUTPUT);
auto size = ShaderArg<uint32_t>(0, SHADER_ARG_OUTPUT);
const struct s expected[] = {
{ 1, 0xfbfcfdff12345678, 0xa112 }
};
run_shader(kernel_source, inout.size(), 1, 1, inout, size);
for (int i = 0; i < inout.size(); ++i) {
EXPECT_EQ(inout[i].uc, expected[i].uc);
EXPECT_EQ(inout[i].ul, expected[i].ul);
EXPECT_EQ(inout[i].us, expected[i].us);
}
EXPECT_EQ(size, sizeof(struct s));
}
TEST_F(ComputeTest, packed_struct_arg)
{
#pragma pack(push, 1)
struct s { uint8_t uc; uint64_t ul; uint16_t us; };
#pragma pack(pop)
const char *kernel_source =
"struct __attribute__((packed)) s {uchar uc; ulong ul; ushort us; };\n\
__kernel void main_test(__global struct s *out, struct s in)\n\
{\n\
uint idx = get_global_id(0);\n\
out[idx].uc = in.uc + 0x12;\n\
out[idx].ul = in.ul + 0x123456789abcdef;\n\
out[idx].us = in.us + 0x1234;\n\
}\n";
auto out = ShaderArg<struct s>({0, 0, 0}, SHADER_ARG_OUTPUT);
auto in = ShaderArg<struct s>({1, 2, 3}, SHADER_ARG_INPUT);
const struct s expected[] = {
{ 0x12 + 1, 0x123456789abcdef + 2, 0x1234 + 3 }
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].uc, expected[i].uc);
EXPECT_EQ(out[i].ul, expected[i].ul);
EXPECT_EQ(out[i].us, expected[i].us);
}
}
TEST_F(ComputeTest, packed_struct_local)
{
#pragma pack(push, 1)
struct s { uint8_t uc; uint64_t ul; uint16_t us; };
#pragma pack(pop)
const char *kernel_source =
"struct __attribute__((packed)) s {uchar uc; ulong ul; ushort us; };\n\
__kernel void main_test(__global struct s *out, __constant struct s *in)\n\
{\n\
uint idx = get_global_id(0);\n\
__local struct s tmp[2];\n\
tmp[get_local_id(0)] = in[idx];\n\
barrier(CLK_LOCAL_MEM_FENCE);\n\
out[idx] = tmp[(get_local_id(0) + 1) % 2];\n\
}\n";
auto out = ShaderArg<struct s>({{0, 0, 0}, {0, 0, 0}}, SHADER_ARG_OUTPUT);
auto in = ShaderArg<struct s>({{1, 2, 3}, {0x12, 0x123456789abcdef, 0x1234} }, SHADER_ARG_INPUT);
const struct s expected[] = {
{ 0x12, 0x123456789abcdef, 0x1234 },
{ 1, 2, 3 },
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].uc, expected[i].uc);
EXPECT_EQ(out[i].ul, expected[i].ul);
EXPECT_EQ(out[i].us, expected[i].us);
}
}
/* DISABLED because current release versions of WARP either return
* rubbish from reads or crash: they are not prepared to handle
* non-float global constants */
TEST_F(ComputeTest, DISABLED_packed_struct_const)
{
#pragma pack(push, 1)
struct s { uint8_t uc; uint64_t ul; uint16_t us; };
#pragma pack(pop)
const char *kernel_source =
"struct __attribute__((packed)) s {uchar uc; ulong ul; ushort us; };\n\
__kernel void main_test(__global struct s *out, struct s in)\n\
{\n\
__constant struct s base[] = {\n\
{0x12, 0x123456789abcdef, 0x1234},\n\
{0x11, 0x123456789abcdee, 0x1233},\n\
};\n\
uint idx = get_global_id(0);\n\
out[idx].uc = base[idx % 2].uc + in.uc;\n\
out[idx].ul = base[idx % 2].ul + in.ul;\n\
out[idx].us = base[idx % 2].us + in.us;\n\
}\n";
auto out = ShaderArg<struct s>(std::vector<struct s>(2, {0, 0, 0}), SHADER_ARG_OUTPUT);
auto in = ShaderArg<struct s>({1, 2, 3}, SHADER_ARG_INPUT);
const struct s expected[] = {
{ 0x12 + 1, 0x123456789abcdef + 2, 0x1234 + 3 },
{ 0x11 + 1, 0x123456789abcdee + 2, 0x1233 + 3 },
};
run_shader(kernel_source, out.size(), 1, 1, out, in);
for (int i = 0; i < out.size(); ++i) {
EXPECT_EQ(out[i].uc, expected[i].uc);
EXPECT_EQ(out[i].ul, expected[i].ul);
EXPECT_EQ(out[i].us, expected[i].us);
}
}
TEST_F(ComputeTest, DISABLED_printf)
{
const char *kernel_source = R"(
__kernel void main_test(__global float *src, __global uint *dest)
{
__constant char *format_str = "%s: %f";
__constant char *str_val = "Test";
*dest = printf(format_str, str_val, src[0]);
})";
auto src = ShaderArg<float>({ 1.0f }, SHADER_ARG_INPUT);
auto dest = ShaderArg<uint32_t>({ 0xdeadbeef }, SHADER_ARG_OUTPUT);
run_shader(kernel_source, 1, 1, 1, src, dest);
EXPECT_EQ(dest[0], 0);
}
TEST_F(ComputeTest, vload_half)
{
const char *kernel_source = R"(
__kernel void main_test(__global half *src, __global float4 *dest)
{
int offset = get_global_id(0);
dest[offset] = vload_half4(offset, src);
})";
auto src = ShaderArg<uint16_t>({ 0x3c00, 0x4000, 0x4200, 0x4400,
0x4500, 0x4600, 0x4700, 0x4800 }, SHADER_ARG_INPUT);
auto dest = ShaderArg<float>({ FLT_MAX, FLT_MAX, FLT_MAX, FLT_MAX,
FLT_MAX, FLT_MAX, FLT_MAX, FLT_MAX }, SHADER_ARG_OUTPUT);
run_shader(kernel_source, 2, 1, 1, src, dest);
for (unsigned i = 0; i < 8; ++i)
EXPECT_FLOAT_EQ(dest[i], (float)(i + 1));
}
TEST_F(ComputeTest, vstore_half)
{
const char *kernel_source = R"(
__kernel void main_test(__global half *dst, __global float4 *src)
{
int offset = get_global_id(0);
vstore_half4(src[offset], offset, dst);
})";
auto dest = ShaderArg<uint16_t>({0xdead, 0xdead, 0xdead, 0xdead,
0xdead, 0xdead, 0xdead, 0xdead}, SHADER_ARG_OUTPUT);
auto src = ShaderArg<float>({ 1.0, 2.0, 3.0, 4.0,
5.0, 6.0, 7.0, 8.0 }, SHADER_ARG_INPUT);
run_shader(kernel_source, 2, 1, 1, dest, src);
const uint16_t expected[] = { 0x3c00, 0x4000, 0x4200, 0x4400,
0x4500, 0x4600, 0x4700, 0x4800 };
for (unsigned i = 0; i < 8; ++i)
EXPECT_EQ(dest[i], expected[i]);
}
TEST_F(ComputeTest, inline_function)
{
const char *kernel_source = R"(
inline float helper(float foo)
{
return foo * 2;
}
__kernel void main_test(__global float *dst, __global float *src)
{
*dst = helper(*src);
})";
auto dest = ShaderArg<float>({ NAN }, SHADER_ARG_OUTPUT);
auto src = ShaderArg<float>({ 1.0f }, SHADER_ARG_INPUT);
run_shader(kernel_source, 1, 1, 1, dest, src);
EXPECT_EQ(dest[0], 2.0f);
}
TEST_F(ComputeTest, unused_arg)
{
const char *kernel_source = R"(
__kernel void main_test(__global int *dst, __global int *unused, __global int *src)
{
int i = get_global_id(0);
dst[i] = src[i];
})";
auto dest = ShaderArg<int>({ -1, -1, -1, -1 }, SHADER_ARG_OUTPUT);
auto src = ShaderArg<int>({ 1, 2, 3, 4 }, SHADER_ARG_INPUT);
auto unused = ShaderArg<int>({ -1, -1, -1, -1 }, SHADER_ARG_INPUT);
run_shader(kernel_source, 4, 1, 1, dest, unused, src);
for (int i = 0; i < 4; ++i)
EXPECT_EQ(dest[i], i + 1);
}
TEST_F(ComputeTest, spec_constant)
{
const char *spirv_asm = R"(
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
%1 = OpExtInstImport "OpenCL.std"
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %2 "main_test" %__spirv_BuiltInGlobalInvocationId
%4 = OpString "kernel_arg_type.main_test.uint*,"
OpSource OpenCL_C 102000
OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId"
OpName %output "output"
OpName %entry "entry"
OpName %output_addr "output.addr"
OpName %id "id"
OpName %call "call"
OpName %conv "conv"
OpName %idxprom "idxprom"
OpName %arrayidx "arrayidx"
OpName %add "add"
OpName %mul "mul"
OpName %idxprom1 "idxprom1"
OpName %arrayidx2 "arrayidx2"
OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId
OpDecorate %__spirv_BuiltInGlobalInvocationId Constant
OpDecorate %id Alignment 4
OpDecorate %output_addr Alignment 8
OpDecorate %uint_1 SpecId 1
%ulong = OpTypeInt 64 0
%uint = OpTypeInt 32 0
%uint_1 = OpSpecConstant %uint 1
%v3ulong = OpTypeVector %ulong 3
%_ptr_Input_v3ulong = OpTypePointer Input %v3ulong
%void = OpTypeVoid
%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint
%24 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint
%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint
%_ptr_Function_uint = OpTypePointer Function %uint
%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3ulong Input
%2 = OpFunction %void DontInline %24
%output = OpFunctionParameter %_ptr_CrossWorkgroup_uint
%entry = OpLabel
%output_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function
%id = OpVariable %_ptr_Function_uint Function
OpStore %output_addr %output Aligned 8
%27 = OpLoad %v3ulong %__spirv_BuiltInGlobalInvocationId Aligned 32
%call = OpCompositeExtract %ulong %27 0
%conv = OpUConvert %uint %call
OpStore %id %conv Aligned 4
%28 = OpLoad %_ptr_CrossWorkgroup_uint %output_addr Aligned 8
%29 = OpLoad %uint %id Aligned 4
%idxprom = OpUConvert %ulong %29
%arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %28 %idxprom
%30 = OpLoad %uint %arrayidx Aligned 4
%31 = OpLoad %uint %id Aligned 4
%add = OpIAdd %uint %31 %uint_1
%mul = OpIMul %uint %30 %add
%32 = OpLoad %_ptr_CrossWorkgroup_uint %output_addr Aligned 8
%33 = OpLoad %uint %id Aligned 4
%idxprom1 = OpUConvert %ulong %33
%arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %32 %idxprom1
OpStore %arrayidx2 %mul Aligned 4
OpReturn
OpFunctionEnd)";
Shader shader = assemble(spirv_asm);
Shader spec_shader = specialize(shader, 1, 5);
auto inout = ShaderArg<uint32_t>({ 0x00000001, 0x10000001, 0x00020002, 0x04010203 },
SHADER_ARG_INOUT);
const uint32_t expected[] = {
0x00000005, 0x60000006, 0x000e000e, 0x20081018
};
CompileArgs args = { inout.size(), 1, 1 };
run_shader(spec_shader, args, inout);
for (int i = 0; i < inout.size(); ++i)
EXPECT_EQ(inout[i], expected[i]);
}