| /external/OpenCL-CTS/test_conformance/device_execution/ |
| D | execute_block.cpp | 32 NL, "int __constant globalVar = 7;"
33 NL, "int (^__constant globalBlock)(int) = ^int(int num)"
34 NL, "{"
35 NL, " return globalVar * num * (1+ get_global_id(0));"
36 NL, "};"
37 NL, "kernel void block_global_scope(__global int* res)"
38 NL, "{"
39 NL, " size_t tid = get_global_id(0);"
40 NL, " res[tid] = -1;"
41 NL, " res[tid] = globalBlock(3) - 21*(tid + 1);"
[all …]
|
| D | enqueue_flags.cpp | 34 NL, "#define BITS_DEPTH " STRINGIFY_VALUE(BITS_DEPTH)
35 NL, ""
36 NL, "void block_fn(__global int* array, int index, size_t ls, size_t gs, __global int* res)"
37 NL, "{"
38 NL, " int val = 0;"
39 NL, " size_t lid = get_local_id(0);"
40 NL, " size_t tid = get_global_id(0);"
41 NL, ""
42 NL, " if(tid == 0)"
43 NL, " {"
[all …]
|
| D | enqueue_block.cpp | 32 NL, "void block_fn(size_t tid, int mul, __global int* res)"
33 NL, "{"
34 NL, " res[tid] = mul * 7 - 21;"
35 NL, "}"
36 NL, ""
37 NL, "kernel void enqueue_simple_block(__global int* res)"
38 NL, "{"
39 NL, " int multiplier = 3;"
40 NL, " size_t tid = get_global_id(0);"
41 NL, ""
[all …]
|
| D | enqueue_wg_size.cpp | 34 NL, "void block_fn(int level, int maxGlobalWorkSize, __global int* rnd, __global int* res)"
35 NL, "{"
36 NL, " size_t tidX = get_global_id(0);"
37 NL, " queue_t def_q = get_default_queue();"
38 NL, " if(--level < 0) return;"
39 NL, ""
40 NL, " void (^kernelBlock)(void) = ^{ block_fn(level, maxGlobalWorkSize, rnd, res); };"
41 NL, " uint wg = get_kernel_work_group_size(kernelBlock);"
42 NL, ""
43 NL, " const size_t gs = 64 * 64 * 64;"
[all …]
|
| D | enqueue_ndrange.cpp | 31 NL,
32 "void block_fn(int len, __global atomic_uint* val)" NL,
33 "{" NL,
35 "memory_order_relaxed, memory_scope_device);" NL,
36 "}" NL,
37 "" NL,
40 "atomic_uint* val, __global uint* ofs_arr)" NL,
41 "{" NL,
42 " size_t tid = get_global_id(0);" NL,
43 " void (^kernelBlock)(void) = ^{ block_fn(len, val); };" NL,
[all …]
|
| D | nested_blocks.cpp | 35 NL, "void block_fn(__global int* res, int level)"
36 NL, "{"
37 NL, " size_t tid = get_global_id(0);"
38 NL, " queue_t def_q = get_default_queue();"
39 NL, " ndrange_t ndrange = ndrange_1D(3);"
40 NL, " if(--level < 0) return;"
41 NL, ""
42 NL, " void (^kernelBlock)(void) = ^{ block_fn(res, level); };"
43 NL, ""
44 NL, " // Only 1 work-item enqueues block"
[all …]
|
| D | host_multi_queue.cpp | 32 NL, "void block_fn(size_t tid, int mul, __global int* res)"
33 NL, "{"
34 NL, " res[tid] = mul * 7 - 21;"
35 NL, "}"
36 NL, ""
37 NL, "kernel void multi_queue_simple_block1(__global int* res)"
38 NL, "{"
39 NL, " int multiplier = 3;"
40 NL, " size_t tid = get_global_id(0);"
41 NL, ""
[all …]
|
| D | enqueue_multi_queue.cpp | 33 NL "#define BLOCK_COMPLETED 0"
34 NL "#define BLOCK_SUBMITTED 1"
35 NL ""
36 NL "kernel void enqueue_block_multi_queue(__global int* res, __global int* buff %s)"
37 NL "{"
38 NL " uint i, n = %d;"
39 NL " clk_event_t block_evt[%d];"
40 NL " queue_t q[] = { %s };"
41 NL " queue_t *queue = q;"
42 NL ""
[all …]
|
| D | host_queue_order.cpp | 33 NL, "void block_fn(uint num, __global int* res)"
34 NL, "{"
35 NL, " size_t tid = get_global_id(0);"
36 NL, ""
37 NL, " for(int i = 1 ; i < tid ; i++)"
38 NL, " {"
39 NL, " for(int j = 0 ; j < num ; j++)"
40 NL, " atomic_add(res+tid, 1);"
41 NL, " }"
42 NL, "}"
[all …]
|
| /external/OpenCL-CTS/test_conformance/generic_address_space/ |
| D | basic_tests.cpp | 99 NL in test_function_get_fence()
100 NL "__global int gint = 1;" in test_function_get_fence()
101 NL "__global uchar guchar = 3;" in test_function_get_fence()
102 NL in test_function_get_fence()
103 … NL "bool helperFunction(int *intp, float *floatp, uchar *ucharp, ushort *ushortp, long *longp) {" in test_function_get_fence()
104 NL " if (!isFenceValid(get_fence(intp)))" in test_function_get_fence()
105 NL " return false;" in test_function_get_fence()
106 NL " if (!isFenceValid(get_fence(floatp)))" in test_function_get_fence()
107 NL " return false;" in test_function_get_fence()
108 NL " if (!isFenceValid(get_fence(ucharp)))" in test_function_get_fence()
[all …]
|
| D | advanced_tests.cpp | 267 NL in test_library_function()
268 NL "bool helperFunction(float *floatp, float val) {" in test_library_function()
269 NL " if (!isFenceValid(get_fence(floatp)))" in test_library_function()
270 NL " return false;" in test_library_function()
271 NL in test_library_function()
272 NL " if (*floatp != val)" in test_library_function()
273 NL " return false;" in test_library_function()
274 NL in test_library_function()
275 NL " return true;" in test_library_function()
276 NL "}" in test_library_function()
[all …]
|
| D | base.h | 25 #define NL "\n" macro
29 NL
30 NL "// current spec says get_fence can return any valid fence"
31 NL "bool isFenceValid(cl_mem_fence_flags fence) {"
32 … NL " if ((fence == 0) || (fence == CLK_GLOBAL_MEM_FENCE) || (fence == CLK_LOCAL_MEM_FENCE) "
33 NL " || (fence == (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE)))"
34 NL " return true;"
35 NL " else"
36 NL " return false;"
37 NL "}"
[all …]
|
| D | stress_tests.cpp | 114 NL in test_max_number_of_params()
115 NL "bool helperFunction(int *ptr0 ", in test_max_number_of_params()
118 NL " // check first pointer only" in test_max_number_of_params()
119 NL " if (!isFenceValid(get_fence(ptr0)))" in test_max_number_of_params()
120 NL " return false;" in test_max_number_of_params()
121 NL in test_max_number_of_params()
122 NL " return true;" in test_max_number_of_params()
123 NL "}" in test_max_number_of_params()
124 NL in test_max_number_of_params()
125 NL "__kernel void testKernel(__global uint *results) {" in test_max_number_of_params()
[all …]
|
| /external/OpenCL-CTS/test_conformance/non_uniform_work_group/ |
| D | TestNonUniformWorkGroup.cpp | 19 #define NL "\n" macro
26 NL "#define MAX_DIMS 3"
27 NL "typedef struct _DataContainerAttrib"
28 NL "{"
29 NL " unsigned long get_global_size[MAX_DIMS];"
30 NL " unsigned long get_global_offset[MAX_DIMS];"
31 NL " unsigned long get_local_size[MAX_DIMS];"
32 NL " unsigned long get_enqueued_local_size[MAX_DIMS];"
33 NL " unsigned long get_global_id[MAX_DIMS];"
34 NL " unsigned long get_local_id[MAX_DIMS];"
[all …]
|
| /external/deqp/external/openglcts/modules/gles31/ |
| D | es31cShaderStorageBufferObjectTests.cpp | 75 const char* const kGLSLVer = "#version 310 es" NL "precision highp float;" NL "precision highp int;…
503 …NL "layout(std430, binding = 1) buffer InputBuffer {" NL " vec4 position[3];" NL "} g_input_buffe… in Run()
504 "void main() {" NL " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}"; in Run()
505 const char* const glsl_fs = NL "layout(location = 0) out vec4 o_color;" NL "void main() {" NL in Run()
506 " o_color = vec4(0.0, 1.0, 0.0, 1.0);" NL "}"; in Run()
557 …const char* const glsl_cs = NL "layout(local_size_x = 1) in;" NL "buffer Buffer {" NL " int resul… in Run()
558 "void main() {" NL " result = 7;" NL "}"; in Run()
906 …NL "layout(std430) buffer Buffer {" NL " vec4 position[3];" NL "} g_input_buffer;" NL "void main(… in Run()
907 " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}", in Run()
908 NL "coherent buffer Buffer {" NL " buffer vec4 position0;" NL " coherent vec4 position1;" NL in Run()
[all …]
|
| D | es31cComputeShaderTests.cpp | 513 …NL "layout(local_size_x = 1, local_size_y = 1) in;" NL "layout(std430) buffer Output {" NL " vec4… in Setup()
514 "} g_out;" NL "void main() {" NL " g_out.data = vec4(1.0, 2.0, 3.0, 4.0);" NL "}"; in Setup()
576 return NL "1. Verify that declared local work size has correct effect." NL in Purpose()
577 "2. Verify that the number of shader invocations is correct." NL in Purpose()
578 …3. Verify that the built-in variables: gl_WorkGroupSize, gl_WorkGroupID, gl_GlobalInvocationID," NL in Purpose()
579 " gl_LocalInvocationID and gl_LocalInvocationIndex has correct values." NL in Purpose()
585 return NL "1. Create several CS with various local sizes." NL in Method()
586 "2. Dispatch each CS with DispatchCompute and DispatchComputeIndirect commands." NL in Method()
602 ss << NL "layout(local_size_x = " << x << ", local_size_y = " << y << ", local_size_z = " << z in GenSource()
603 << ") in;" NL "layout(std430, binding = " << binding in GenSource()
[all …]
|
| D | es31cShaderImageLoadStoreTests.cpp | 55 const char* const kGLSLSIA = NL "#extension GL_OES_shader_image_atomic : require";
58 NL "precision highp float;" NL "precision highp int;" NL "precision highp sampler2D;" NL
59 …"precision highp sampler3D;" NL "precision highp samplerCube;" NL "precision highp sampler2DArray;…
60 …"precision highp isampler2D;" NL "precision highp isampler3D;" NL "precision highp isamplerCube;" …
61 …"precision highp isampler2DArray;" NL "precision highp usampler2D;" NL "precision highp usampler3D…
62 …"precision highp usamplerCube;" NL "precision highp usampler2DArray;" NL "precision highp image2D;…
63 "precision highp image3D;" NL "precision highp imageCube;" NL "precision highp image2DArray;" NL
64 "precision highp iimage2D;" NL "precision highp iimage3D;" NL "precision highp iimageCube;" NL
65 … "precision highp iimage2DArray;" NL "precision highp uimage2D;" NL "precision highp uimage3D;" NL
66 "precision highp uimageCube;" NL "precision highp uimage2DArray;";
[all …]
|
| D | es31cShaderAtomicCountersTests.cpp | 51 return NL ""; in Title()
55 return NL ""; in Purpose()
59 return NL ""; in Method()
63 return NL ""; in PassCriteria()
790 return NL "Atomic Counters usage in the Compute Shader stage"; in Title()
794 return NL "Verify that atomic counters work as expected in the Compute Shader stage." NL in Purpose()
795 "In particular make sure that values returned by GLSL built-in functions" NL in Purpose()
796 "atomicCounterIncrement and atomicCounterDecrement are unique in every shader invocation." NL in Purpose()
801 return NL ""; in Method()
805 return NL ""; in PassCriteria()
[all …]
|
| D | es31cShaderImageSizeTests.cpp | 46 …"#version 310 es\n" NL "precision highp float;" NL "precision highp int;" NL "precision highp imag…
47 "precision highp image3D;" NL "precision highp imageCube;" NL "precision highp image2DArray;" NL
48 "precision highp iimage2D;" NL "precision highp iimage3D;" NL "precision highp iimageCube;" NL
49 "precision highp iimage2DArray;" NL "precision highp uimage2D;" NL "precision highp uimage3D;" NL
50 "precision highp uimageCube;" NL "precision highp uimage2DArray;";
57 return NL ""; in Title()
61 return NL ""; in Purpose()
65 return NL ""; in Method()
69 return NL ""; in PassCriteria()
392 os << NL "#define KSIZE 4"; in GenShader()
[all …]
|
| /external/deqp/external/openglcts/modules/gl/ |
| D | gl4cShaderStorageBufferObjectTests.cpp | 656 …NL "layout(std430, binding = 1) buffer InputBuffer {" NL " vec4 position[3];" NL "} g_input_buffe… in Run()
657 "void main() {" NL " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}"; in Run()
659 const char* const glsl_fs = NL "layout(location = 0) out vec4 o_color;" NL "void main() {" NL in Run()
660 " o_color = vec4(0.0, 1.0, 0.0, 1.0);" NL "}"; in Run()
710 …const char* const glsl_cs = NL "layout(local_size_x = 1) in;" NL "buffer Buffer {" NL " int resul… in Run()
711 "void main() {" NL " result = 7;" NL "}"; in Run()
1080 …NL "layout(std430) buffer Buffer {" NL " vec4 position[3];" NL "} g_input_buffer;" NL "void main(… in Run()
1081 " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}", in Run()
1082 NL "coherent buffer Buffer {" NL " buffer vec4 position0;" NL " coherent vec4 position1;" NL in Run()
1083 " restrict readonly vec4 position2;" NL "} g_input_buffer;" NL "void main() {" NL in Run()
[all …]
|
| D | gl4cES31CompatibilityShaderStorageBufferObjectTests.cpp | 77 const char* const kGLSLVer = "#version 310 es" NL "precision highp float;" NL "precision highp int;…
506 …NL "layout(std430, binding = 1) buffer InputBuffer {" NL " vec4 position[3];" NL "} g_input_buffe… in Run()
507 "void main() {" NL " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}"; in Run()
508 const char* const glsl_fs = NL "layout(location = 0) out vec4 o_color;" NL "void main() {" NL in Run()
509 " o_color = vec4(0.0, 1.0, 0.0, 1.0);" NL "}"; in Run()
560 …const char* const glsl_cs = NL "layout(local_size_x = 1) in;" NL "buffer Buffer {" NL " int resul… in Run()
561 "void main() {" NL " result = 7;" NL "}"; in Run()
905 …NL "layout(std430) buffer Buffer {" NL " vec4 position[3];" NL "} g_input_buffer;" NL "void main(… in Run()
906 " gl_Position = g_input_buffer.position[gl_VertexID];" NL "}", in Run()
907 NL "coherent buffer Buffer {" NL " buffer vec4 position0;" NL " coherent vec4 position1;" NL in Run()
[all …]
|
| D | gl4cComputeShaderTests.cpp | 504 …NL "layout(local_size_x = 1, local_size_y = 1) in;" NL "layout(std430) buffer Output {" NL " vec4… in Setup()
505 "} g_out;" NL "void main() {" NL " g_out.data = vec4(1.0, 2.0, 3.0, 4.0);" NL "}"; in Setup()
567 return NL "1. Verify that declared local work size has correct effect." NL in Purpose()
568 "2. Verify that the number of shader invocations is correct." NL in Purpose()
569 …3. Verify that the built-in variables: gl_WorkGroupSize, gl_WorkGroupID, gl_GlobalInvocationID," NL in Purpose()
570 " gl_LocalInvocationID and gl_LocalInvocationIndex has correct values." NL in Purpose()
576 return NL "1. Create several CS with various local sizes." NL in Method()
577 "2. Dispatch each CS with DispatchCompute and DispatchComputeIndirect commands." NL in Method()
593 ss << NL "layout(local_size_x = " << x << ", local_size_y = " << y << ", local_size_z = " << z in GenSource()
594 << ") in;" NL "layout(std430, binding = " << binding in GenSource()
[all …]
|
| D | gl4cES31CompatibilityShaderImageLoadStoreTests.cpp | 56 const char* const kGLSLSIA = NL "#extension GL_OES_shader_image_atomic : require";
58 NL "precision highp float;" NL "precision highp int;" NL "precision highp sampler2D;" NL
59 …"precision highp sampler3D;" NL "precision highp samplerCube;" NL "precision highp sampler2DArray;…
60 …"precision highp isampler2D;" NL "precision highp isampler3D;" NL "precision highp isamplerCube;" …
61 …"precision highp isampler2DArray;" NL "precision highp usampler2D;" NL "precision highp usampler3D…
62 …"precision highp usamplerCube;" NL "precision highp usampler2DArray;" NL "precision highp image2D;…
63 "precision highp image3D;" NL "precision highp imageCube;" NL "precision highp image2DArray;" NL
64 "precision highp iimage2D;" NL "precision highp iimage3D;" NL "precision highp iimageCube;" NL
65 … "precision highp iimage2DArray;" NL "precision highp uimage2D;" NL "precision highp uimage3D;" NL
66 "precision highp uimageCube;" NL "precision highp uimage2DArray;";
[all …]
|
| D | gl4cShaderImageLoadStoreTests.cpp | 1410 …const char* src_vs = "#version 420 core" NL "layout(location = 0) in vec4 i_position;" NL "void ma… in Write()
1411 " gl_Position = i_position;" NL "}"; in Write()
1476 …os << "#version 420 core" NL "layout(" << FormatEnumToString(internalformat) << ") writeonly unifo… in GenFS()
1477 …<< TypePrefix<T>() << "image2D g_image;" NL "void main() {" NL " ivec2 coord = ivec2(gl_FragCoord… in GenFS()
1479 << TypePrefix<T>() << "vec4" << value << ");" NL " discard;" NL "}"; in GenFS()
1599 …const char* src_vs = "#version 420 core" NL "layout(location = 0) in vec4 i_position;" NL "void ma… in Read()
1600 " gl_Position = i_position;" NL "}"; in Read()
1658 os << "#version 420 core" NL "layout(location = 0) out vec4 o_color;" NL "layout(" in GenFS()
1660 << "image2D g_image;" NL "void main() {" NL " ivec2 coord = ivec2(gl_FragCoord.xy);" NL " " in GenFS()
1661 …<< TypePrefix<T>() << "vec4 v = imageLoad(g_image, coord);" NL " if (v != " << TypePrefix<T>() <<… in GenFS()
[all …]
|
| D | gl4cShaderAtomicCountersTests.cpp | 52 return NL ""; in Title()
57 return NL ""; in Purpose()
62 return NL ""; in Method()
67 return NL ""; in PassCriteria()
894 return NL "Atomic Counter Buffer - basic operations"; in Title()
899 return NL in Purpose()
900 "Verify that basic buffer operations work as expected with new buffer target." NL in Purpose()
901 …ted commands: BindBuffer, BufferData, BufferSubData, MapBuffer, MapBufferRange, UnmapBuffer and" NL in Purpose()
907 return NL ""; in Method()
912 return NL ""; in PassCriteria()
[all …]
|