Lines Matching defs:inout
46 __kernel void main_test(__global struct shift *inout)\n\
51 lc[id] = inout[id];\n\
52 inout[id2].ret = (ushort) lc[id2].val << (ushort) lc[id2].shift;\n\
55 auto inout = ShaderArg<struct shift>({
63 run_shader(kernel_source, inout.size(), 1, 1, inout);
64 for (int i = 0; i < inout.size(); ++i)
65 EXPECT_EQ(inout[i].ret, expected[i]);
325 "__kernel void main_test(__global unsigned char *inout)\n\
328 inout[idx] = inout[idx] + 1;\n\
330 auto inout = ShaderArg<uint8_t> ({ 100, 110, 120, 130 }, SHADER_ARG_INOUT);
334 run_shader(kernel_source, inout.size(), 1, 1, inout);
335 for (int i = 0; i < inout.size(); ++i)
336 EXPECT_EQ(inout[i], expected[i]);
342 "__kernel void main_test(__global unsigned short *inout)\n\
345 inout[idx] = inout[idx] + 1;\n\
347 auto inout = ShaderArg<uint16_t> ({ 10000, 10010, 10020, 10030 }, SHADER_ARG_INOUT);
351 run_shader(kernel_source, inout.size(), 1, 1, inout);
352 for (int i = 0; i < inout.size(); ++i)
353 EXPECT_EQ(inout[i], expected[i]);
359 "__kernel void main_test(__global unsigned long *inout)\n\
362 inout[idx] = inout[idx] + 1;\n\
365 auto inout = ShaderArg<uint64_t>({ base, base + 10, base + 20, base + 30 },
370 run_shader(kernel_source, inout.size(), 1, 1, inout);
371 for (int i = 0; i < inout.size(); ++i)
372 EXPECT_EQ(inout[i], expected[i]);
401 auto inout = ShaderArg<uint32_t>({0x00000001, 0x10000001, 0x00020002, 0x04010203},
406 run_shader(kernel_source, inout.size(), 1, 1, inout);
407 for (int i = 0; i < inout.size(); ++i)
408 EXPECT_EQ(inout[i], expected[i]);
547 "__kernel void main_test(__global ulong *inout)\n\
556 inout[idx] = tmp[idx];\n\
558 auto inout = ShaderArg<uint64_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
562 run_shader(kernel_source, inout.size(), 1, 1, inout);
563 for (int i = 0; i < inout.size(); ++i)
564 EXPECT_EQ(inout[i], expected[i]);
570 "__kernel void main_test(__global ushort *inout)\n\
579 inout[idx] = tmp[idx];\n\
581 auto inout = ShaderArg<uint16_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
585 run_shader(kernel_source, inout.size(), 1, 1, inout);
586 for (int i = 0; i < inout.size(); ++i)
587 EXPECT_EQ(inout[i], expected[i]);
594 __kernel void main_test(__global uint *inout)\n\
604 inout[idx] = mul + trunc(tmp[idx].f[1]);\n\
606 auto inout = ShaderArg<uint32_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
608 run_shader(kernel_source, inout.size(), 1, 1, inout);
609 for (int i = 0; i < inout.size(); ++i)
610 EXPECT_EQ(inout[i], expected[i]);
616 "__kernel void main_test(__global uint *inout)\n\
625 inout[idx] = tmp[idx];\n\
627 auto inout = ShaderArg<uint32_t>({ 0, 0, 0, 0 }, SHADER_ARG_INOUT);
631 run_shader(kernel_source, inout.size(), 1, 1, inout);
632 for (int i = 0; i < inout.size(); ++i)
633 EXPECT_EQ(inout[i], expected[i]);
647 auto inout = ShaderArg<struct two_vals>({ { 8, 8 }, { 16, 16 }, { 64, 64 }, { 65536, 65536 } },
655 run_shader(kernel_source, inout.size(), 1, 1, inout);
656 for (int i = 0; i < inout.size(); ++i) {
657 EXPECT_EQ(inout[i].add, expected[i].add);
658 EXPECT_EQ(inout[i].mul, expected[i].mul);
666 "__kernel void main_test(__global uint2 *inout)\n\
669 inout[id].x = inout[id].x + id;\n\
670 inout[id].y = inout[id].y * id;\n\
672 auto inout = ShaderArg<struct uint2>({ { 8, 8 }, { 16, 16 }, { 64, 64 }, { 65536, 65536 } },
680 run_shader(kernel_source, inout.size(), 1, 1, inout);
681 for (int i = 0; i < inout.size(); ++i) {
682 EXPECT_EQ(inout[i].x, expected[i].x);
683 EXPECT_EQ(inout[i].y, expected[i].y);
691 "__kernel void main_test(__global ushort2 *inout)\n\
694 inout[id].x = inout[id].x + id;\n\
695 inout[id].y = inout[id].y * id;\n\
697 auto inout = ShaderArg<struct ushort2>({ { 8, 8 }, { 16, 16 }, { 64, 64 },
706 run_shader(kernel_source, inout.size(), 1, 1, inout);
707 for (int i = 0; i < inout.size(); ++i) {
708 EXPECT_EQ(inout[i].x, expected[i].x);
709 EXPECT_EQ(inout[i].y, expected[i].y);
717 "__kernel void main_test(__global uchar3 *inout)\n\
720 inout[id].x = inout[id].x + id;\n\
721 inout[id].y = inout[id].y * id;\n\
722 inout[id].z = inout[id].y + inout[id].x;\n\
724 auto inout = ShaderArg<struct uchar3>({ { 8, 8, 8 }, { 16, 16, 16 }, { 64, 64, 64 }, { 255, 255, 255 } },
732 run_shader(kernel_source, inout.size(), 1, 1, inout);
733 for (int i = 0; i < inout.size(); ++i) {
734 EXPECT_EQ(inout[i].x, expected[i].x);
735 EXPECT_EQ(inout[i].y, expected[i].y);
736 EXPECT_EQ(inout[i].z, expected[i].z);
776 "__kernel void main_test(__global uint8 *inout)\n\
779 inout[id].s01234567 = inout[id].s01234567 * 2;\n\
781 auto inout = ShaderArg<struct uint8>({ { 1, 2, 3, 4, 5, 6, 7, 8 } },
786 run_shader(kernel_source, inout.size(), 1, 1, inout);
787 for (int i = 0; i < inout.size(); ++i) {
788 EXPECT_EQ(inout[i].s0, expected[i].s0);
789 EXPECT_EQ(inout[i].s1, expected[i].s1);
790 EXPECT_EQ(inout[i].s2, expected[i].s2);
791 EXPECT_EQ(inout[i].s3, expected[i].s3);
792 EXPECT_EQ(inout[i].s4, expected[i].s4);
793 EXPECT_EQ(inout[i].s5, expected[i].s5);
794 EXPECT_EQ(inout[i].s6, expected[i].s6);
795 EXPECT_EQ(inout[i].s7, expected[i].s7);
805 R"(__kernel void main_test(__global ulong16 *inout)
809 local_array[id] = inout[id];
811 inout[id] = local_array[0] * 2;
813 auto inout = ShaderArg<struct ulong16>({ { 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15 } },
818 run_shader(kernel_source, inout.size(), 1, 1, inout);
819 for (int i = 0; i < inout.size(); ++i) {
821 EXPECT_EQ(inout[i].values[j], expected[i].values[j]);
924 "__kernel void main_test(__global int *inout)\n\
926 inout[get_global_id(0)] = inout[get_global_id(0)] % 3;\n\
928 auto inout = ShaderArg<int32_t>({ -4, -3, -2, -1, 0, 1, 2, 3, 4 },
933 run_shader(kernel_source, inout.size(), 1, 1, inout);
934 for (int i = 0; i < inout.size(); ++i)
935 EXPECT_EQ(inout[i], expected[i]);
941 "__kernel void main_test(__global int *inout)\n\
943 inout[get_global_id(0)] = inout[get_global_id(0)] % -3;\n\
945 auto inout = ShaderArg<int32_t>({ -4, -3, -2, -1, 0, 1, 2, 3, 4 },
950 run_shader(kernel_source, inout.size(), 1, 1, inout);
951 for (int i = 0; i < inout.size(); ++i)
952 EXPECT_EQ(inout[i], expected[i]);
958 "__kernel void main_test(__global uint *inout)\n\
960 inout[get_global_id(0)] = inout[get_global_id(0)] % 0xfffffffc;\n\
962 auto inout = ShaderArg<uint32_t>({ 0xfffffffa, 0xfffffffb, 0xfffffffc, 0xfffffffd, 0xfffffffe },
967 run_shader(kernel_source, inout.size(), 1, 1, inout);
968 for (int i = 0; i < inout.size(); ++i)
969 EXPECT_EQ(inout[i], expected[i]);
975 "__kernel void main_test(__global uint *inout)\n\
977 inout[get_global_id(0)] = rotate(inout[get_global_id(0)], (uint)get_global_id(0) * 4);\n\
979 auto inout = ShaderArg<uint32_t>(std::vector<uint32_t>(4, 0xdeadbeef),
984 run_shader(kernel_source, inout.size(), 1, 1, inout);
985 for (int i = 0; i < inout.size(); ++i)
986 EXPECT_EQ(inout[i], expected[i]);
992 "__kernel void main_test(__global uint *inout)\n\
994 inout[get_global_id(0)] = popcount(inout[get_global_id(0)]);\n\
996 auto inout = ShaderArg<uint32_t>({ 0, 0x1, 0x3, 0x101, 0x110011, ~0u },
1001 run_shader(kernel_source, inout.size(), 1, 1, inout);
1002 for (int i = 0; i < inout.size(); ++i)
1003 EXPECT_EQ(inout[i], expected[i]);
1009 "__kernel void main_test(__global uint *inout)\n\
1011 inout[get_global_id(0)] = hadd(inout[get_global_id(0)], 1u << 31);\n\
1013 auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, 0xfffffffc, 0xfffffffd, 0xfffffffe, 0xffffffff },
1025 run_shader(kernel_source, inout.size(), 1, 1, inout);
1026 for (int i = 0; i < inout.size(); ++i)
1027 EXPECT_EQ(inout[i], expected[i]);
1033 "__kernel void main_test(__global uint *inout)\n\
1035 inout[get_global_id(0)] = rhadd(inout[get_global_id(0)], 1u << 31);\n\
1037 auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, 0xfffffffc, 0xfffffffd, 0xfffffffe, 0xffffffff },
1049 run_shader(kernel_source, inout.size(), 1, 1, inout);
1050 for (int i = 0; i < inout.size(); ++i)
1051 EXPECT_EQ(inout[i], expected[i]);
1057 "__kernel void main_test(__global uint *inout)\n\
1059 inout[get_global_id(0)] = add_sat(inout[get_global_id(0)], 2u);\n\
1061 auto inout = ShaderArg<uint32_t>({ 0xffffffff - 3, 0xffffffff - 2, 0xffffffff - 1, 0xffffffff },
1066 run_shader(kernel_source, inout.size(), 1, 1, inout);
1067 for (int i = 0; i < inout.size(); ++i)
1068 EXPECT_EQ(inout[i], expected[i]);
1074 "__kernel void main_test(__global uint *inout)\n\
1076 inout[get_global_id(0)] = sub_sat(inout[get_global_id(0)], 2u);\n\
1078 auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3 }, SHADER_ARG_INOUT);
1082 run_shader(kernel_source, inout.size(), 1, 1, inout);
1083 for (int i = 0; i < inout.size(); ++i)
1084 EXPECT_EQ(inout[i], expected[i]);
1090 "__kernel void main_test(__global uint *inout)\n\
1092 inout[get_global_id(0)] = mul_hi(inout[get_global_id(0)], 1u << 31);\n\
1094 auto inout = ShaderArg<uint32_t>({ 0, 1, 2, 3, (1u << 31) }, SHADER_ARG_INOUT);
1098 run_shader(kernel_source, inout.size(), 1, 1, inout);
1099 for (int i = 0; i < inout.size(); ++i)
1100 EXPECT_EQ(inout[i], expected[i]);
1106 "__kernel void main_test(__global float *inout)\n\
1108 inout[get_global_id(0)] = ldexp(inout[get_global_id(0)], 5);\n\
1110 auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 2.0f }, SHADER_ARG_INOUT);
1114 run_shader(kernel_source, inout.size(), 1, 1, inout);
1115 for (int i = 0; i < inout.size(); ++i)
1116 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1122 "__kernel void main_test(__global float *inout)\n\
1124 inout[get_global_id(0)] = ldexp(inout[get_global_id(0)], get_global_id(0));\n\
1126 auto inout = ShaderArg<float>({ 0.25f, 0.5f, 0.75f, 1.0f }, SHADER_ARG_INOUT);
1130 run_shader(kernel_source, inout.size(), 1, 1, inout);
1131 for (int i = 0; i < inout.size(); ++i)
1132 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1138 "__kernel void main_test(__global float *inout)\n\
1141 inout[get_global_id(0)] = frexp(inout[get_global_id(0)], &exp);\n\
1143 auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 3.0f }, SHADER_ARG_INOUT);
1147 run_shader(kernel_source, inout.size(), 1, 1, inout);
1148 for (int i = 0; i < inout.size(); ++i)
1149 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1155 "__kernel void main_test(__global float *inout)\n\
1158 frexp(inout[get_global_id(0)], &exp);\n\
1159 inout[get_global_id(0)] = (float)exp;\n\
1161 auto inout = ShaderArg<float>({ 0.0f, 0.5f, 1.0f, 3.0f }, SHADER_ARG_INOUT);
1165 run_shader(kernel_source, inout.size(), 1, 1, inout);
1166 for (int i = 0; i < inout.size(); ++i)
1167 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1173 "__kernel void main_test(__global uint *inout)\n\
1175 inout[get_global_id(0)] = clz(inout[get_global_id(0)]);\n\
1177 auto inout = ShaderArg<uint32_t>({ 0, 1, 0xffff, (1u << 30), (1u << 31) }, SHADER_ARG_INOUT);
1181 run_shader(kernel_source, inout.size(), 1, 1, inout);
1182 for (int i = 0; i < inout.size(); ++i)
1183 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1191 __kernel void main_test(__global struct sin_vals *inout)\n\
1193 inout[get_global_id(0)].clc = sin(inout[get_global_id(0)].in);\n\
1194 inout[get_global_id(0)].native = native_sin(inout[get_global_id(0)].in);\n\
1202 auto inout = ShaderArg<sin_vals>(input, SHADER_ARG_INOUT);
1209 run_shader(kernel_source, inout.size(), 1, 1, inout);
1210 for (int i = 0; i < inout.size(); ++i) {
1211 EXPECT_FLOAT_EQ(inout[i].in, inout[i].in);
1212 EXPECT_FLOAT_EQ(inout[i].clc, inout[i].clc);
1213 EXPECT_NEAR(inout[i].clc, inout[i].native, 0.008f); // range from DXIL spec
1220 "__kernel void main_test(__global float *inout)\n\
1222 inout[get_global_id(0)] = cosh(inout[get_global_id(0)]);\n\
1224 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1228 run_shader(kernel_source, inout.size(), 1, 1, inout);
1229 for (int i = 0; i < inout.size(); ++i)
1230 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1236 "__kernel void main_test(__global float *inout)\n\
1238 inout[get_global_id(0)] = native_exp(inout[get_global_id(0)]);\n\
1240 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1244 run_shader(kernel_source, inout.size(), 1, 1, inout);
1245 for (int i = 0; i < inout.size(); ++i)
1246 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1252 "__kernel void main_test(__global float *inout)\n\
1254 inout[get_global_id(0)] = native_exp10(inout[get_global_id(0)]);\n\
1256 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1260 run_shader(kernel_source, inout.size(), 1, 1, inout);
1261 for (int i = 0; i < inout.size(); ++i)
1262 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1268 "__kernel void main_test(__global float *inout)\n\
1270 inout[get_global_id(0)] = native_exp2(inout[get_global_id(0)]);\n\
1272 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1276 run_shader(kernel_source, inout.size(), 1, 1, inout);
1277 for (int i = 0; i < inout.size(); ++i)
1278 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1284 "__kernel void main_test(__global float *inout)\n\
1286 inout[get_global_id(0)] = native_log(inout[get_global_id(0)]);\n\
1288 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1292 run_shader(kernel_source, inout.size(), 1, 1, inout);
1293 for (int i = 0; i < inout.size(); ++i)
1294 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1300 "__kernel void main_test(__global float *inout)\n\
1302 inout[get_global_id(0)] = native_log10(inout[get_global_id(0)]);\n\
1304 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1308 run_shader(kernel_source, inout.size(), 1, 1, inout);
1309 for (int i = 0; i < inout.size(); ++i)
1310 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1316 "__kernel void main_test(__global float *inout)\n\
1318 inout[get_global_id(0)] = native_log2(inout[get_global_id(0)]);\n\
1320 auto inout = ShaderArg<float>({ 0.0f, 1.0f, 2.0f, 3.0f }, SHADER_ARG_INOUT);
1324 run_shader(kernel_source, inout.size(), 1, 1, inout);
1325 for (int i = 0; i < inout.size(); ++i)
1326 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1332 "__kernel void main_test(__global float *inout)\n\
1334 inout[get_global_id(0)] = rint(inout[get_global_id(0)]);\n\
1337 auto inout = ShaderArg<float>({ 0.5f, 1.5f, -0.5f, -1.5f, 1.4f }, SHADER_ARG_INOUT);
1341 run_shader(kernel_source, inout.size(), 1, 1, inout);
1342 for (int i = 0; i < inout.size(); ++i)
1343 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1349 "__kernel void main_test(__global float *inout)\n\
1351 inout[get_global_id(0)] = round(inout[get_global_id(0)]);\n\
1353 auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
1358 run_shader(kernel_source, inout.size(), 1, 1, inout);
1359 for (int i = 0; i < inout.size(); ++i)
1360 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1366 "__kernel void main_test(__global float *inout, float mul)\n\
1368 inout[get_global_id(0)] = inout[get_global_id(0)] * mul;\n\
1370 auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
1376 run_shader(kernel_source, inout.size(), 1, 1, inout, mul);
1377 for (int i = 0; i < inout.size(); ++i)
1378 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1410 __kernel void main_test(__global float *inout)\n\
1412 inout[get_global_id(0)] = foo(inout[get_global_id(0)]);\n\
1415 auto inout = ShaderArg<float>({ 2.0f }, SHADER_ARG_INOUT);
1419 run_shader(srcs, inout.size(), 1, 1, inout);
1420 for (int i = 0; i < inout.size(); ++i)
1421 EXPECT_EQ(inout[i], expected[i]);
1439 __kernel void main_test(__global float *inout)\n\
1441 inout[get_global_id(0)] = foo(inout[get_global_id(0)]);\n\
1448 auto inout = ShaderArg<float>({ 2.0f }, SHADER_ARG_INOUT);
1452 run_shader(exe, { (unsigned)inout.size(), 1, 1 }, inout);
1453 for (int i = 0; i < inout.size(); ++i)
1454 EXPECT_EQ(inout[i], expected[i]);
1461 void main_test(__global float *inout)\n\
1464 tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
1465 tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
1467 inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
1470 auto inout = ShaderArg<float>({ 2.0f, 4.0f }, SHADER_ARG_INOUT);
1474 run_shader(kernel_source, inout.size(), 1, 1, inout);
1475 for (int i = 0; i < inout.size(); ++i)
1476 EXPECT_EQ(inout[i], expected[i]);
1483 __kernel void main_test(__global uchar *inout)\n\
1486 tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
1487 tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
1489 inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
1492 auto inout = ShaderArg<uint8_t>({ 2, 4 }, SHADER_ARG_INOUT);
1494 run_shader(kernel_source, inout.size(), 1, 1, inout);
1495 for (int i = 0; i < inout.size(); ++i)
1496 EXPECT_EQ(inout[i], expected[i]);
1633 "__kernel void main_test(__global uint *inout, __local uint2 *tmp)\n\
1635 tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
1636 tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
1638 inout[get_global_id(0)] = tmp[get_local_id(0) % 2].x * tmp[(get_local_id(0) + 1) % 2].y;\n\
1640 auto inout = ShaderArg<uint32_t>({ 2, 4 }, SHADER_ARG_INOUT);
1643 run_shader(kernel_source, inout.size(), 1, 1, inout, tmp);
1644 for (int i = 0; i < inout.size(); ++i)
1645 EXPECT_EQ(inout[i], expected[i]);
1652 "__kernel void main_test(__global uint *inout, __local uint2 *tmp, __local uint *tmp2)\n\
1654 tmp[get_local_id(0)].x = inout[get_global_id(0)] + 1;\n\
1655 tmp[get_local_id(0)].y = inout[get_global_id(0)] - 1;\n\
1658 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\
1660 auto inout = ShaderArg<uint32_t>({ 2, 4 }, SHADER_ARG_INOUT);
1664 run_shader(kernel_source, inout.size(), 1, 1, inout, tmp, tmp2);
1665 for (int i = 0; i < inout.size(); ++i)
1666 EXPECT_EQ(inout[i], expected[i]);
1688 "__kernel __attribute__((vec_type_hint(float4))) void main_test(__global float *inout)\n\
1690 inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
1700 "__kernel __attribute__((vec_type_hint(uchar2))) void main_test(__global float *inout)\n\
1702 inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
1712 "__kernel void main_test(__global float *inout)\n\
1714 inout[get_global_id(0)] *= inout[get_global_id(1)];\n\
1726 "__kernel void main_test(__global float *inout, float mul)\n\
1728 inout[get_global_id(0)] = inout[get_global_id(0)] * mul;\n\
1730 auto inout = ShaderArg<float>({ 0, 0.3f, -0.3f, 0.5f, -0.5f, 1.1f, -1.1f },
1744 run_shader(kernel_source, inout.size(), 1, 1, inout, mul);
1745 for (int i = 0; i < inout.size(); ++i)
1746 EXPECT_FLOAT_EQ(inout[i], expected[i]);
1769 "__kernel void main_test(__global int *inout, __global int *old)\n\
1771 old[get_global_id(0)] = atomic_add(inout + get_global_id(0), 3);\n\
1773 auto inout = ShaderArg<int32_t>({ 2, 4 }, SHADER_ARG_INOUT);
1777 run_shader(kernel_source, inout.size(), 1, 1, inout, old);
1778 for (int i = 0; i < inout.size(); ++i) {
1779 EXPECT_EQ(inout[i], expected_inout[i]);
1787 "__kernel void main_test(__global int *inout, __global int *old)\n\
1789 old[get_global_id(0)] = atomic_min(inout + get_global_id(0), 1);\n\
1791 auto inout = ShaderArg<int32_t>({ 0, 2, -1 }, SHADER_ARG_INOUT);
1795 run_shader(kernel_source, inout.size(), 1, 1, inout, old);
1796 for (int i = 0; i < inout.size(); ++i) {
1797 EXPECT_EQ(inout[i], expected_inout[i]);
1806 __kernel void main_test(__global int *inout)\n\
1808 atomic_and(inout, ~(1 << get_global_id(0)));\n\
1809 atomic_or(inout, (1 << (get_global_id(0) + 4)));\n\
1811 auto inout = ShaderArg<int32_t>(0xf, SHADER_ARG_INOUT);
1813 run_shader(kernel_source, 3, 1, 1, inout);
1814 for (int i = 0; i < inout.size(); ++i)
1815 EXPECT_EQ(inout[i], expected[i]);
1822 __kernel void main_test(__global int *inout)\n\
1824 while (atomic_cmpxchg(inout, get_global_id(0), get_global_id(0) + 1) != get_global_id(0))\n\
1827 auto inout = ShaderArg<int32_t>(0, SHADER_ARG_INOUT);
1829 run_shader(kernel_source, 2, 1, 1, inout);
1830 for (int i = 0; i < inout.size(); ++i)
1831 EXPECT_EQ(inout[i], expected_inout[i]);
1838 __kernel void main_test(__global ushort *inout)\n\
1842 atomic_or(&tmp, inout[get_global_id(0)] << (get_global_id(0) * 8));\n\
1844 inout[get_global_id(0)] = tmp;\n\
1846 auto inout = ShaderArg<uint16_t>({ 2, 4 }, SHADER_ARG_INOUT);
1848 run_shader(kernel_source, inout.size(), 1, 1, inout);
1849 for (int i = 0; i < inout.size(); ++i)
1850 EXPECT_EQ(inout[i], expected[i]);
2032 __kernel void main_test(__global struct s *inout, global uint *size)\n\
2035 inout[idx].uc = idx + 1;\n\
2036 inout[idx].ul = ((ulong)(idx + 1 + 0xfbfcfdfe) << 32) | 0x12345678;\n\
2037 inout[idx].us = ((ulong)(idx + 1 + 0xa0) << 8) | 0x12;\n\
2040 auto inout = ShaderArg<struct s>({0, 0, 0}, SHADER_ARG_OUTPUT);
2046 run_shader(kernel_source, inout.size(), 1, 1, inout, size);
2047 for (int i = 0; i < inout.size(); ++i) {
2048 EXPECT_EQ(inout[i].uc, expected[i].uc);
2049 EXPECT_EQ(inout[i].ul, expected[i].ul);
2050 EXPECT_EQ(inout[i].us, expected[i].us);
2301 auto inout = ShaderArg<uint32_t>({ 0x00000001, 0x10000001, 0x00020002, 0x04010203 },
2306 CompileArgs args = { (unsigned)inout.size(), 1, 1 };
2307 run_shader(spec_shader, args, inout);
2308 for (int i = 0; i < inout.size(); ++i)
2309 EXPECT_EQ(inout[i], expected[i]);