diff --git a/benchmarks/opencl/ma/compact-features.cl b/benchmarks/opencl/ma/compact-features.cl new file mode 100644 index 0000000000..b8f37d8a04 --- /dev/null +++ b/benchmarks/opencl/ma/compact-features.cl @@ -0,0 +1,41 @@ +// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable \ +// -finline-functions -finline-hint-functions \ +// -emit-llvm -fno-discard-value-names -c compact-features.cl +// -o compact-features.bc -DNUM_VAR=1 -DNUM_OUTPUT_VAR=1 +// llvm-spirv compact-features.bc -o compact-features.spv +// spirv-dis compact-features.spv > compact-features.spvasm + +#ifdef DV2WG +#define scope memory_scope_work_group +#else +#define scope memory_scope_device +#endif + +#ifdef LC2GB +#define bar_flag CLK_GLOBAL_MEM_FENCE +#else +#define bar_flag CLK_LOCAL_MEM_FENCE +#endif + +__kernel void compact_features(__global uint* flags, + __global uint* out_indices, + __global uint* group_offset) { + __local uint s_idx; + + uint tid = get_local_id(0); + uint gid = get_global_id(0); + uint group_id = get_group_id(0); + + // The work-group leader initializes the index using a plain store. + if (tid == 0) { + s_idx = group_offset[group_id]; + } + + barrier(bar_flag); + + // Threads filter data and contend for slots in the output list. + if (flags[gid]) { + uint dst = atomic_fetch_add_explicit((atomic_uint*)&s_idx, 1, memory_order_relaxed, scope); + out_indices[dst] = gid; + } +} \ No newline at end of file diff --git a/benchmarks/opencl/ma/histogram-implicit.cl b/benchmarks/opencl/ma/histogram-implicit.cl new file mode 100644 index 0000000000..be87185285 --- /dev/null +++ b/benchmarks/opencl/ma/histogram-implicit.cl @@ -0,0 +1,33 @@ +// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable \ +// -finline-functions -finline-hint-functions \ +// -emit-llvm -fno-discard-value-names -c histogram.cl +// -o histogram.bc -DNUM_VAR=1 -DNUM_OUTPUT_VAR=1 +// llvm-spirv histogram.bc -o histogram.spv +// spirv-dis histogram.spv > histogram.spvasm + +#define HIST_BINS 2 + +__kernel void histo_main_kernel(global uint *sm_mappings, + global uint *global_histo) +{ + __local uint sub_histo[HIST_BINS]; + + int tid = get_local_id(0); + int gid = get_global_id(0); + + // Safe plain store because threads own distinct indices; + sub_histo[tid] = 0; + + barrier(CLK_LOCAL_MEM_FENCE); + + // Multiple threads contend for the same bins; + uint bin_index = sm_mappings[gid]; + atom_add(sub_histo + bin_index, 1); + + barrier(CLK_LOCAL_MEM_FENCE); + + // Read local result plain to flush to global; + uint count = sub_histo[tid]; + if (count > 0) + atom_add(global_histo + tid, count); +} \ No newline at end of file diff --git a/benchmarks/opencl/ma/histogram.cl b/benchmarks/opencl/ma/histogram.cl new file mode 100644 index 0000000000..cdcdf12af8 --- /dev/null +++ b/benchmarks/opencl/ma/histogram.cl @@ -0,0 +1,45 @@ +// clang -x cl -cl-std=CL2.0 -target spir-unknown-unknown -cl-opt-disable \ +// -finline-functions -finline-hint-functions \ +// -emit-llvm -fno-discard-value-names -c histogram.cl +// -o histogram.bc -DNUM_VAR=1 -DNUM_OUTPUT_VAR=1 +// llvm-spirv histogram.bc -o histogram.spv +// spirv-dis histogram.spv > histogram.spvasm + +#ifdef DV2WG +#define scope memory_scope_work_group +#else +#define scope memory_scope_device +#endif + +#ifdef LC2GB +#define flag CLK_GLOBAL_MEM_FENCE +#else +#define flag CLK_LOCAL_MEM_FENCE +#endif + +#define HIST_BINS 2 + +__kernel void histo_main_kernel(__global uint *sm_mappings, + __global uint *global_histo) +{ + __local uint sub_histo[HIST_BINS]; + + int tid = get_local_id(0); + int gid = get_global_id(0); + + // Safe plain store because threads own distinct indices; + sub_histo[tid] = 0; + + barrier(CLK_LOCAL_MEM_FENCE); + + // Multiple threads contend for the same bins; + uint bin_index = sm_mappings[gid]; + atomic_fetch_add_explicit((atomic_uint*)&sub_histo[bin_index], 1, memory_order_relaxed, memory_scope_work_group); + + barrier(CLK_LOCAL_MEM_FENCE); + + // Read local result plain to flush to global; + uint count = sub_histo[tid]; + if (count > 0) + atomic_fetch_add_explicit((atomic_uint*)&global_histo[tid], count, memory_order_relaxed, scope); +} \ No newline at end of file diff --git a/benchmarks/opencl/mixed-atomicity/histogram.cl b/benchmarks/opencl/mixed-atomicity/histogram.cl new file mode 100644 index 0000000000..c9f438644f --- /dev/null +++ b/benchmarks/opencl/mixed-atomicity/histogram.cl @@ -0,0 +1,33 @@ +#ifdef DV2WG +#define scope memory_scope_work_group +#else +#define scope memory_scope_device +#endif + +#ifdef LC2GB +#define flag CLK_GLOBAL_MEM_FENCE +#else +#define flag CLK_LOCAL_MEM_FENCE +#endif + +__kernel void histo_main_kernel(__global uint *sm_mappings, __global uint *global_histo) +{ + __local uint sub_histo[2]; + + int tid = get_local_id(0); + int gid = get_global_id(0); + + sub_histo[tid] = 0; + + barrier(flag); + + uint bin_index = sm_mappings[gid]; + + atomic_fetch_add_explicit((atomic_uint*)&sub_histo[bin_index], 1, memory_order_relaxed, memory_scope_work_group); + + barrier(flag); + + uint count = sub_histo[tid]; + if (count > 0) + atomic_fetch_add_explicit((atomic_uint*)&global_histo[tid], count, memory_order_relaxed, scope); +} \ No newline at end of file diff --git a/cat/opencl-ma.cat b/cat/opencl-ma.cat new file mode 100644 index 0000000000..66074f16d0 --- /dev/null +++ b/cat/opencl-ma.cat @@ -0,0 +1,129 @@ +OpenCL +(* OpenCL Memory Model *) + +(* +* This model is based on: +* https://multicore.doc.ic.ac.uk/overhauling/opencl_base.cat +* https://multicore.doc.ic.ac.uk/overhauling/opencl_scopedsc.cat +*) + +// Base relations: +// wi: same work-item (same thread) +// swg: same work-group +// sdv: same device +// syncbar: same barrier id + +// Tags: +// WI: work-item scope +// WG: work-group scope +// DV: device scope +// ALL: all-svm-devices scope +// GLOBAL: global memory +// LOCAL: local memory + +// dynamic_tag relates events to itself that access an address whose init event is marked X or Fence tagged with X +let dynamic_tag(X) = [range([IW & X]; loc)] | [X & F] + +let mo = co \ (NA * NA) +let sb = po +let rb = rf^-1;mo | ([R] \ [range(rf)]);loc;[W] +let unv = _ * _ +let wi = int + +(* Inclusive scopes *) +let incl = (swg & (WG * WG)) | (sdv & (DV * DV)) | (ALL * ALL) + +(*******************) +(* Synchronisation *) +(*******************) + +let Acq = (ACQ | SC | ACQ_REL) & (R | F) +let Rel = (REL | SC | ACQ_REL) & (W | F) + +(* Fences sequenced before or after *) +let Fsb = [F]; sb +let sbF = sb; [F] + +(* Release sequence *) +let rs_prime = (_ * RMW) | wi +let rs = mo & rs_prime & ~((mo & ~rs_prime) ; mo) + +(* Release-acquire synchronisation *) +let ra_sw(r) = ((r & [Rel]); Fsb?; [W & A]; rs?; r; rf; [R & A]; sbF?; ([Acq] & r)) & incl & ~wi + +(* Barrier synchronisation *) +// in OpenCL a barrier results in two fence operations: entry and exit fences: +// https://registry.khronos.org/OpenCL/specs/3.0-unified/html/OpenCL_API.html#_work_group_functions +// In our implementation, we use a single barrier event and omit the special EF and XF tags +let bar_sw(r) = r; syncbar & ~wi & swg; r + +(* Allowed to synchronise on the other region *) +let scf = (SC * SC) | ((dynamic_tag(GLOBAL) & dynamic_tag(LOCAL)); unv; (dynamic_tag(GLOBAL) & dynamic_tag(LOCAL))) + +(* Global and local synchronises-with *) +let gsw = ra_sw(dynamic_tag(GLOBAL)) | bar_sw(dynamic_tag(GLOBAL)) | (scf & ra_sw(dynamic_tag(LOCAL))) +let lsw = ra_sw(dynamic_tag(LOCAL)) | bar_sw(dynamic_tag(LOCAL)) | (scf & ra_sw(dynamic_tag(GLOBAL))) + +(******************) +(* Happens-before *) +(******************) + +(* Global and local happens-before *) +// Since we use single barrier events, we exclude the identity relation from the happens-before relation +let ini = IW * ~IW +let ghb = ((dynamic_tag(GLOBAL); sb; dynamic_tag(GLOBAL)) | (dynamic_tag(GLOBAL); ini; dynamic_tag(GLOBAL)) | gsw)+ \ id +let lhb = ((dynamic_tag(LOCAL); sb; dynamic_tag(LOCAL)) | (dynamic_tag(LOCAL); ini; dynamic_tag(LOCAL)) | lsw)+ \ id + +irreflexive ghb as O-HbG +irreflexive lhb as O-HbL + +(*************) +(* Coherence *) +(*************) + +let coh(hb) = (rf^-1)?; mo; rf?; hb +irreflexive coh(ghb) as O-CohG +irreflexive coh(lhb) as O-CohL + +(************************) +(* Consistency of reads *) +(************************) + +(* A load can only read from a store that already happened. *) +irreflexive rf; (ghb | lhb) as O-Rf + +(* Visible side effects *) +let vis(hb) = (W * R) & hb & loc & ~((hb & loc); [W]; hb) + +(* A non-atomic load can only read from a store that is visible. *) +empty (rf; ([NA] & dynamic_tag(GLOBAL))) \ vis(ghb) as O-NaRfG +empty (rf; ([NA] & dynamic_tag(LOCAL))) \ vis(lhb) as O-NaRfL + +(* Consistency of RMWs *) +// The original model was tested with Herd, which treats RMW as a single atomic operation. +// We have modified the model to handle RMW as a sequence of atomic operations. +// irreflexive rf | (mo;mo;rf^-1) | (mo;rf) as O-Rmw +empty rmw & (fre; coe) as 0-Atomic + +(****************************************) +(* Sequential consistency, simplified, *) +(* with scoped SC axioms *) +(****************************************) + +let scp = Fsb?; (rb | mo | (ghb | lhb)); sbF? +acyclic (SC*SC) & scp & incl as O-Sscoped + +(***************) +(* Races *) +(***************) + +(* data_races *) +let cnf = ((W * W) | (W * R) | (R * W)) & loc +// TODO: "there is exactly one initial event per location", in current implementation, memory object like global scope id is initialized by three write events +let dr = cnf & ~(ghb | lhb) & ~(ghb | lhb)^-1 & ~wi & ~incl \ ((_ * IW) | (IW * _)) +flag ~empty dr as data_race + +(* unsequenced_races *) +let symm(r) = r | r^-1 +let ur = (wi & cnf & ~symm(sb)) \ id +flag ~empty ur as unsequenced_race \ No newline at end of file diff --git a/cat/opencl.cat b/cat/opencl.cat index 384f8c2cc5..741309bde3 100644 --- a/cat/opencl.cat +++ b/cat/opencl.cat @@ -50,7 +50,7 @@ let rs_prime = (_ * RMW) | wi let rs = mo & rs_prime & ~((mo & ~rs_prime) ; mo) (* Release-acquire synchronisation *) -let ra_sw(r) = ((r & [Rel]); Fsb?; [W \ WI]; rs?; r; rf; [R \ WI]; sbF?; ([Acq] & r)) & incl & ~wi +let ra_sw(r) = ((r & [Rel]); Fsb?; [W & A]; rs?; r; rf; [R & A]; sbF?; ([Acq] & r)) & incl & ~wi (* Barrier synchronisation *) // in OpenCL a barrier results in two fence operations: entry and exit fences: diff --git a/dartagnan/src/main/antlr4/Spirv.g4 b/dartagnan/src/main/antlr4/Spirv.g4 index 0dca81f3d0..c47bd0e18d 100644 --- a/dartagnan/src/main/antlr4/Spirv.g4 +++ b/dartagnan/src/main/antlr4/Spirv.g4 @@ -1741,6 +1741,7 @@ clspvReflection | clspvReflection_specConstantSubgroupMaxSize | clspvReflection_specConstantWorkDim | clspvReflection_specConstantWorkgroupSize + | clspvReflection_workgroupVariableSize ; clspvReflection_kernel : ModeExt_Kernel kernelIdRef nameIdRef (numArguments (flags attributes?)?)?; @@ -1755,6 +1756,7 @@ clspvReflection_argumentStorageImage : ModeExt_ArgumentStorageImage decl ordinal clspvReflection_argumentSampler : ModeExt_ArgumentSampler decl ordinal descriptorSetIdRef binding argInfo?; clspvReflection_argumentWorkgroup : ModeExt_ArgumentWorkgroup decl ordinal specId elemSize argInfo?; clspvReflection_specConstantWorkgroupSize : ModeExt_SpecConstantWorkgroupSize x y z; +clspvReflection_workgroupVariableSize : ModeExt_WorkgroupVariableSize x y; clspvReflection_specConstantGlobalOffset : ModeExt_SpecConstantGlobalOffset x y z; clspvReflection_specConstantWorkDim : ModeExt_SpecConstantWorkDim dimIdRef; clspvReflection_pushConstantGlobalOffset : ModeExt_PushConstantGlobalOffset offsetIdRef sizeIdRef; diff --git a/dartagnan/src/main/antlr4/SpirvLexer.g4 b/dartagnan/src/main/antlr4/SpirvLexer.g4 index f700372ca7..1616786b1f 100644 --- a/dartagnan/src/main/antlr4/SpirvLexer.g4 +++ b/dartagnan/src/main/antlr4/SpirvLexer.g4 @@ -1836,6 +1836,7 @@ ModeExt_PrintfInfo : 'PrintfInfo'; ModeExt_PrintfBufferStorageBuffer : 'PrintfBufferStorageBuffer'; ModeExt_PrintfBufferPointerPushConstant : 'PrintfBufferPointerPushConstant'; ModeExt_NormalizedSamplerMaskPushConstant : 'NormalizedSamplerMaskPushConstant'; +ModeExt_WorkgroupVariableSize : 'WorkgroupVariableSize'; ModeExt_Round : 'Round'; ModeExt_RoundEven : 'RoundEven'; diff --git a/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/VisitorOpsMemory.java b/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/VisitorOpsMemory.java index 92ae31b6bd..09f149e516 100644 --- a/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/VisitorOpsMemory.java +++ b/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/VisitorOpsMemory.java @@ -47,6 +47,7 @@ public VisitorOpsMemory(ProgramBuilder builder) { @Override public Event visitOpStore(SpirvParser.OpStoreContext ctx) { Expression pointer = builder.getExpression(ctx.pointer().getText()); + pointer.getMemoryObjects().forEach(mo -> mo.addFeatureTag(Tag.C11.NON_ATOMIC_LOCATION)); String valueId = ctx.object().getText(); Expression value = builder.getExpression(valueId); Type type = value.getType(); @@ -63,6 +64,7 @@ public Event visitOpStore(SpirvParser.OpStoreContext ctx) { public Event visitOpLoad(SpirvParser.OpLoadContext ctx) { String resultId = ctx.idResult().getText(); Expression pointer = builder.getExpression(ctx.pointer().getText()); + pointer.getMemoryObjects().forEach(mo -> mo.addFeatureTag(Tag.C11.NON_ATOMIC_LOCATION)); Type type = builder.getType(ctx.idResultType().getText()); List events = visitMemoryAccess(resultId, type, pointer, (i, exp) -> { String regId = resultId; diff --git a/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/extenstions/VisitorExtensionClspvReflection.java b/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/extenstions/VisitorExtensionClspvReflection.java index c25ce2c83c..6e1a30b9a5 100644 --- a/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/extenstions/VisitorExtensionClspvReflection.java +++ b/dartagnan/src/main/java/com/dat3m/dartagnan/parsers/program/visitors/spirv/extenstions/VisitorExtensionClspvReflection.java @@ -58,6 +58,12 @@ public Expression visitClspvReflection_specConstantWorkgroupSize(SpirvParser.Cls return null; } + @Override + public Expression visitClspvReflection_workgroupVariableSize(SpirvParser.ClspvReflection_workgroupVariableSizeContext ctx) { + // Do nothing, will be overwritten by BuiltIn WorkgroupSize + return null; + } + @Override public Expression visitClspvReflection_pushConstantGlobalOffset(SpirvParser.ClspvReflection_pushConstantGlobalOffsetContext ctx) { return setPushConstantValue("PushConstantGlobalOffset", ctx.offsetIdRef().getText(), ctx.sizeIdRef().getText()); @@ -195,7 +201,8 @@ public Set getSupportedInstructions() { "PushConstantNumWorkgroups", "PushConstantRegionOffset", "PushConstantRegionGroupOffset", - "SpecConstantWorkgroupSize" + "SpecConstantWorkgroupSize", + "WorkgroupVariableSize" ); } } diff --git a/dartagnan/src/main/java/com/dat3m/dartagnan/program/event/Tag.java b/dartagnan/src/main/java/com/dat3m/dartagnan/program/event/Tag.java index 3f194e8dff..2d4da89610 100644 --- a/dartagnan/src/main/java/com/dat3m/dartagnan/program/event/Tag.java +++ b/dartagnan/src/main/java/com/dat3m/dartagnan/program/event/Tag.java @@ -378,6 +378,7 @@ public static final class OpenCL { public static final String GLOBAL_SPACE = "GLOBAL"; public static final String LOCAL_SPACE = "LOCAL"; public static final String GENERIC_SPACE = "GENERIC"; + public static final String PRIVATE = "PRIVATE"; // Default Tags public static final String DEFAULT_SPACE = GENERIC_SPACE; public static final String DEFAULT_SCOPE = DEVICE; @@ -563,7 +564,8 @@ public static String toOpenCLTag(String tag) { // Storage class case SC_GENERIC -> OpenCL.GENERIC_SPACE; case SC_FUNCTION, - SC_INPUT, + SC_PRIVATE -> OpenCL.PRIVATE; + case SC_INPUT, SC_WORKGROUP -> OpenCL.LOCAL_SPACE; case SC_UNIFORM_CONSTANT, SC_PHYS_STORAGE_BUFFER, @@ -571,8 +573,7 @@ public static String toOpenCLTag(String tag) { case SC_PUSH_CONSTANT, SC_UNIFORM, SC_OUTPUT, - SC_STORAGE_BUFFER, - SC_PRIVATE -> throw new UnsupportedOperationException( + SC_STORAGE_BUFFER -> throw new UnsupportedOperationException( getErrorMsg(model, "storage class", tag)); default -> throw new IllegalArgumentException( diff --git a/dartagnan/src/main/java/com/dat3m/dartagnan/program/processing/transformers/MemoryTransformer.java b/dartagnan/src/main/java/com/dat3m/dartagnan/program/processing/transformers/MemoryTransformer.java index 9b39db624e..c9a5a825d8 100644 --- a/dartagnan/src/main/java/com/dat3m/dartagnan/program/processing/transformers/MemoryTransformer.java +++ b/dartagnan/src/main/java/com/dat3m/dartagnan/program/processing/transformers/MemoryTransformer.java @@ -125,6 +125,9 @@ private Expression applyMapping(MemoryObject memObj, int scopeDepth) { ? program.getMemory().allocateVirtual(memObj.getKnownSize(), true, null) : program.getMemory().allocate(memObj.getKnownSize()); copy.setName(makeVariableName(scopeDepth, memObj.getName())); + for (String tag : memObj.getFeatureTags()) { + copy.addFeatureTag(tag); + } for (int offset : memObj.getInitializedFields()) { Expression value = memObj.getInitialValue(offset); if (value instanceof NonDetValue) { diff --git a/dartagnan/src/main/resources/log4j2.xml b/dartagnan/src/main/resources/log4j2.xml index d8e1f0f65f..78421c2dc1 100644 --- a/dartagnan/src/main/resources/log4j2.xml +++ b/dartagnan/src/main/resources/log4j2.xml @@ -16,7 +16,7 @@ - + diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMARacesTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMARacesTest.java new file mode 100644 index 0000000000..543aa63ae7 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMARacesTest.java @@ -0,0 +1,41 @@ +package com.dat3m.dartagnan.litmus; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.configuration.Property; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.utils.rules.Provider; +import com.dat3m.dartagnan.utils.rules.Providers; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; + +import java.io.IOException; +import java.util.EnumSet; + +@RunWith(Parameterized.class) +public class LitmusOpenClMARacesTest extends AbstractLitmusTest { + + @Parameterized.Parameters(name = "{index}: {0}, {1}") + public static Iterable data() throws IOException { + return buildLitmusTests("litmus/OPENCL/", "OPENCL", "-MA-DR"); + } + + @Override + protected Provider getTargetProvider() { + return () -> Arch.OPENCL; + } + + @Override + protected Provider> getPropertyProvider() { + return Provider.fromSupplier(() -> EnumSet.of(Property.CAT_SPEC)); + } + + @Override + protected Provider getWmmProvider() { + return Providers.createWmmFromName(() -> "opencl-ma"); + } + + public LitmusOpenClMARacesTest(String path, Result expected) { + super(path, expected); + } +} diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMATest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMATest.java new file mode 100644 index 0000000000..29b9030c51 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/litmus/LitmusOpenClMATest.java @@ -0,0 +1,34 @@ +package com.dat3m.dartagnan.litmus; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.utils.rules.Provider; +import com.dat3m.dartagnan.utils.rules.Providers; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; + +import java.io.IOException; + +@RunWith(Parameterized.class) +public class LitmusOpenClMATest extends AbstractLitmusTest { + + @Parameterized.Parameters(name = "{index}: {0}, {1}") + public static Iterable data() throws IOException { + return buildLitmusTests("litmus/OPENCL/", "OPENCL", "-MA"); + } + + @Override + protected Provider getTargetProvider() { + return () -> Arch.OPENCL; + } + + @Override + protected Provider getWmmProvider() { + return Providers.createWmmFromName(() -> "opencl-ma"); + } + + public LitmusOpenClMATest(String path, Result expected) { + super(path, expected); + } +} diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/others/program/processing/compilation/VisitorSpirvOpenCLTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/others/program/processing/compilation/VisitorSpirvOpenCLTest.java index 13aaa85060..05eaeacaa8 100644 --- a/dartagnan/src/test/java/com/dat3m/dartagnan/others/program/processing/compilation/VisitorSpirvOpenCLTest.java +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/others/program/processing/compilation/VisitorSpirvOpenCLTest.java @@ -49,7 +49,7 @@ public void testLoad() { Set.of( Tag.C11.MO_RELAXED, Tag.OpenCL.DEVICE, Tag.OpenCL.GLOBAL_SPACE) ); doTestLoad( - Set.of(Tag.Spirv.WORKGROUP, Tag.Spirv.ACQUIRE, Tag.Spirv.SEM_WORKGROUP, Tag.Spirv.SC_FUNCTION), + Set.of(Tag.Spirv.WORKGROUP, Tag.Spirv.ACQUIRE, Tag.Spirv.SEM_WORKGROUP, Tag.Spirv.SC_WORKGROUP), Set.of(Tag.OpenCL.WORK_GROUP, Tag.C11.MO_ACQUIRE, Tag.OpenCL.LOCAL_SPACE) ); } @@ -86,7 +86,7 @@ public void testStore() { Set.of( Tag.C11.MO_RELAXED, Tag.OpenCL.DEVICE, Tag.OpenCL.GLOBAL_SPACE) ); doTestStore( - Set.of(Tag.Spirv.WORKGROUP, Tag.Spirv.RELEASE, Tag.Spirv.SEM_WORKGROUP, Tag.Spirv.SC_FUNCTION), + Set.of(Tag.Spirv.WORKGROUP, Tag.Spirv.RELEASE, Tag.Spirv.SEM_WORKGROUP, Tag.Spirv.SC_WORKGROUP), Set.of(Tag.OpenCL.WORK_GROUP, Tag.C11.MO_RELEASE, Tag.OpenCL.LOCAL_SPACE) ); } @@ -127,7 +127,7 @@ public void testSpirvLoad() { Set.of(Tag.C11.MO_ACQUIRE, Tag.OpenCL.WORK_ITEM, Tag.OpenCL.GENERIC_SPACE) ); doTestSpirvLoad( - Set.of(Tag.Spirv.RELAXED, Tag.Spirv.DEVICE, Tag.Spirv.SC_FUNCTION), + Set.of(Tag.Spirv.RELAXED, Tag.Spirv.DEVICE, Tag.Spirv.SC_WORKGROUP), Set.of(Tag.C11.MO_RELAXED, Tag.OpenCL.DEVICE, Tag.OpenCL.LOCAL_SPACE) ); } @@ -169,7 +169,7 @@ public void testSpirvStore() { Set.of(Tag.C11.MO_RELEASE, Tag.OpenCL.WORK_ITEM, Tag.OpenCL.GENERIC_SPACE) ); doTestSpirvStore( - Set.of(Tag.Spirv.RELAXED, Tag.Spirv.DEVICE, Tag.Spirv.SC_FUNCTION), + Set.of(Tag.Spirv.RELAXED, Tag.Spirv.DEVICE, Tag.Spirv.SC_WORKGROUP), Set.of(Tag.C11.MO_RELAXED, Tag.OpenCL.DEVICE, Tag.OpenCL.LOCAL_SPACE) ); } diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvAssertionsTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvAssertionsTest.java new file mode 100644 index 0000000000..1ec1c48d08 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvAssertionsTest.java @@ -0,0 +1,168 @@ +package com.dat3m.dartagnan.spirv.opencl.cav26; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.PROGRAM_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.*; +import static com.dat3m.dartagnan.utils.Result.UNKNOWN; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvAssertionsTest { + + private final String modelPath = getRootPath("cat/opencl-ma.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvAssertionsTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/opencl/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"ma/histogram-1.1.4.spvasm", 2, FAIL}, + {"ma/histogram-2.1.2.spvasm", 2, PASS}, + {"ma/histogram-4.1.1.spvasm", 2, PASS}, + {"ma/histogram-implicit-1.1.4.spvasm", 2, FAIL}, + {"ma/histogram-implicit-4.1.1.spvasm", 2, PASS}, + {"ma/histogram-lc2gb-1.spvasm", 2, FAIL}, + {"ma/histogram-lc2gb-2.spvasm", 2, FAIL}, + {"ma/compact-features-2.1.2.spvasm", 2, PASS}, + {"ma/compact-features-lc2gb.spvasm", 2, FAIL}, + {"alignment/alignment1-array-global.spvasm", 9, PASS}, + {"alignment/alignment1-array-local.spvasm", 9, PASS}, + {"alignment/alignment1-array-pointer.spvasm", 9, PASS}, + {"alignment/alignment1-struct-global.spvasm", 9, PASS}, + {"alignment/alignment1-struct-local.spvasm", 9, PASS}, + {"alignment/alignment1-struct-pointer.spvasm", 9, PASS}, + {"alignment/alignment2-struct-global.spvasm", 17, PASS}, + {"alignment/alignment2-struct-local.spvasm", 17, PASS}, + {"alignment/alignment2-struct-pointer.spvasm", 17, PASS}, + {"alignment/alignment3-struct-global.spvasm", 9, PASS}, + {"alignment/alignment3-struct-local.spvasm", 9, PASS}, + {"alignment/alignment3-struct-pointer.spvasm", 9, PASS}, + {"alignment/alignment4-struct-global.spvasm", 25, PASS}, + {"alignment/alignment4-struct-local.spvasm", 25, PASS}, + {"alignment/alignment4-struct-pointer.spvasm", 25, PASS}, + {"alignment/alignment5-struct-global.spvasm", 17, PASS}, + {"alignment/alignment5-struct-local.spvasm", 17, PASS}, + {"alignment/alignment5-struct-pointer.spvasm", 17, PASS}, + {"barrier/inlining/barrier-inlining-1-forall-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-1-exists-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-1-forall-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-1-exists-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-no-inlining-1-forall-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-no-inlining-1-exists-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-no-inlining-1-forall-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-no-inlining-1-exists-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-2-forall-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-2-exists-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-2-forall-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-2-exists-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-3-forall-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-3-exists-correct.spvasm", 1, PASS}, + {"barrier/inlining/barrier-inlining-3-forall-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-3-exists-wrong.spvasm", 1, FAIL}, + {"barrier/inlining/barrier-inlining-4-forall-correct.spvasm", 3, PASS}, + {"barrier/inlining/barrier-inlining-4-exists-correct.spvasm", 3, PASS}, + {"barrier/inlining/barrier-inlining-4-forall-wrong.spvasm", 3, FAIL}, + {"barrier/inlining/barrier-inlining-4-exists-wrong.spvasm", 3, FAIL}, + {"barrier/inlining/barrier-inlining-5-forall-correct.spvasm", 3, PASS}, + {"barrier/inlining/barrier-inlining-5-exists-correct.spvasm", 3, PASS}, + {"barrier/inlining/barrier-inlining-5-forall-wrong.spvasm", 3, FAIL}, + {"barrier/inlining/barrier-inlining-5-exists-wrong.spvasm", 3, FAIL}, + {"barrier/scope/barrier-inscope-wg.spvasm", 1, PASS}, + {"barrier/scope/barrier-not-inscope-wg.spvasm", 1, FAIL}, + {"basic/idx-overflow.spvasm", 1, PASS}, + {"benchmarks/caslock-1.1.2.spvasm", 2, UNKNOWN}, + {"benchmarks/caslock-2.1.1.spvasm", 2, UNKNOWN}, + {"benchmarks/caslock-acq2rx.spvasm", 2, FAIL}, + {"benchmarks/caslock-rel2rx.spvasm", 2, FAIL}, + {"benchmarks/caslock-dv2wg-2.1.1.spvasm", 2, UNKNOWN}, + {"benchmarks/caslock-dv2wg-1.1.2.spvasm", 2, FAIL}, + {"benchmarks/ticketlock-1.1.2.spvasm", 1, PASS}, + {"benchmarks/ticketlock-2.1.1.spvasm", 1, PASS}, + {"benchmarks/ticketlock-acq2rx.spvasm", 1, FAIL}, + {"benchmarks/ticketlock-rel2rx.spvasm", 1, FAIL}, + {"benchmarks/ticketlock-dv2wg-2.1.1.spvasm", 2, PASS}, + {"benchmarks/ticketlock-dv2wg-1.1.2.spvasm", 1, FAIL}, + {"benchmarks/ttaslock-1.1.2.spvasm", 2, PASS}, + {"benchmarks/ttaslock-2.1.1.spvasm", 2, PASS}, + {"benchmarks/ttaslock-acq2rx.spvasm", 1, FAIL}, + {"benchmarks/ttaslock-rel2rx.spvasm", 1, FAIL}, + {"benchmarks/ttaslock-dv2wg-2.1.1.spvasm", 2, PASS}, + {"benchmarks/ttaslock-dv2wg-1.1.2.spvasm", 1, FAIL}, + {"benchmarks/xf-barrier-2.1.2.spvasm", 9, PASS}, + {"benchmarks/xf-barrier-2.1.1.spvasm", 9, PASS}, + {"benchmarks/xf-barrier-fail1.spvasm", 9, FAIL}, + {"benchmarks/xf-barrier-fail2.spvasm", 9, FAIL}, + {"benchmarks/xf-barrier-fail3.spvasm", 9, FAIL}, + {"benchmarks/xf-barrier-fail4.spvasm", 9, FAIL}, + {"benchmarks/xf-barrier-weakest.spvasm", 9, FAIL}, + {"patterns/corr.spvasm", 2, PASS}, + {"patterns/iriw.spvasm", 2, PASS}, + {"patterns/mp.spvasm", 2, PASS}, + {"patterns/mp-acq2rx.spvasm", 2, FAIL}, + {"patterns/mp-rel2rx.spvasm", 2, FAIL}, + {"patterns/sb.spvasm", 2, PASS}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.Z3); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.OPENCL); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(PROGRAM_SPEC)); + } +} diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvRacesTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvRacesTest.java new file mode 100644 index 0000000000..b2c9f5f54d --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/cav26/SpirvRacesTest.java @@ -0,0 +1,94 @@ +package com.dat3m.dartagnan.spirv.opencl.cav26; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.CAT_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.FAIL; +import static com.dat3m.dartagnan.utils.Result.PASS; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvRacesTest { + + private final String modelPath = getRootPath("cat/opencl-ma.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvRacesTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/opencl/ma/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"histogram-1.1.4.spvasm", 2, PASS}, + {"histogram-2.1.2.spvasm", 2, PASS}, + {"histogram-4.1.1.spvasm", 2, PASS}, + {"histogram-dv2wg.spvasm", 2, FAIL}, + {"histogram-lc2gb-1.spvasm", 2, FAIL}, + {"histogram-lc2gb-2.spvasm", 2, FAIL}, + {"histogram-implicit-1.1.4.spvasm", 2, PASS}, + {"histogram-implicit-2.1.2.spvasm", 2, FAIL}, + {"histogram-implicit-4.1.1.spvasm", 2, PASS}, + {"compact-features-2.1.2.spvasm", 2, PASS}, + {"compact-features-lc2gb.spvasm", 2, FAIL}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.YICES2); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.OPENCL); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(CAT_SPEC)); + } +} \ No newline at end of file diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/gpuverify/SpirvRacesTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/gpuverify/SpirvRacesTest.java index 303fa30f11..5604849b60 100644 --- a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/gpuverify/SpirvRacesTest.java +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/gpuverify/SpirvRacesTest.java @@ -113,6 +113,7 @@ public static Iterable data() throws IOException { {"atomics/refined_atomic_abstraction/intra_local_counters.spvasm", 1, PASS}, // Should pass according to gpu-verify, suspecting a bug in the memory model + // Accessing NAL using ATOMIC instructions {"atomics/counter.spvasm", 1, FAIL}, // In gpu-verify fails barrier divergence but not leading to a data race @@ -127,7 +128,7 @@ public static Iterable data() throws IOException { // barrier avvis variations {"inter_group_and_barrier_flag_tests/fail/missing_local_barrier_flag.spvasm", 1, FAIL}, - {"inter_group_and_barrier_flag_tests/pass/local_barrier_flag.spvasm", 1, FAIL}, + {"inter_group_and_barrier_flag_tests/pass/local_barrier_flag.spvasm", 1, PASS}, // Unsupported large array (4K elements) leading to OOM // {"misc/fail/2d_array_race.spvasm", 1, FAIL}, diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvAssertionsTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvAssertionsTest.java new file mode 100644 index 0000000000..6deaafeb55 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvAssertionsTest.java @@ -0,0 +1,93 @@ +package com.dat3m.dartagnan.spirv.opencl.ma; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.PROGRAM_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.*; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvAssertionsTest { + + private final String modelPath = getRootPath("cat/opencl-ma.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvAssertionsTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/opencl/ma/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"histogram-1.1.4.spvasm", 2, FAIL}, + {"histogram-2.1.2.spvasm", 2, PASS}, + {"histogram-4.1.1.spvasm", 2, PASS}, + {"histogram-implicit-1.1.4.spvasm", 2, FAIL}, + {"histogram-implicit-4.1.1.spvasm", 2, PASS}, + {"histogram-implicit-2.1.2.spvasm", 2, PASS}, + {"histogram-lc2gb-1.spvasm", 2, FAIL}, + {"histogram-lc2gb-2.spvasm", 2, FAIL}, + {"compact-features-2.1.2.spvasm", 2, PASS}, + {"compact-features-lc2gb.spvasm", 2, FAIL}, + {"compact-features-implicit-2.1.2.spvasm", 2, PASS}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.Z3); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.OPENCL); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(PROGRAM_SPEC)); + } +} diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvRacesTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvRacesTest.java new file mode 100644 index 0000000000..1d634c9b6f --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/opencl/ma/SpirvRacesTest.java @@ -0,0 +1,94 @@ +package com.dat3m.dartagnan.spirv.opencl.ma; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.CAT_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.*; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvRacesTest { + + private final String modelPath = getRootPath("cat/opencl-ma.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvRacesTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/opencl/ma/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"histogram-1.1.4.spvasm", 2, PASS}, + {"histogram-2.1.2.spvasm", 2, PASS}, + {"histogram-4.1.1.spvasm", 2, PASS}, + {"histogram-dv2wg.spvasm", 2, FAIL}, + {"histogram-lc2gb-1.spvasm", 2, FAIL}, + {"histogram-lc2gb-2.spvasm", 2, FAIL}, + {"compact-features-2.1.2.spvasm", 2, PASS}, + {"compact-features-lc2gb.spvasm", 2, FAIL}, + {"histogram-implicit-1.1.4.spvasm", 2, PASS}, + {"histogram-implicit-2.1.2.spvasm", 2, FAIL}, + {"histogram-implicit-4.1.1.spvasm", 2, PASS}, + {"compact-features-implicit-2.1.2.spvasm", 2, PASS}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.YICES2); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.OPENCL); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(CAT_SPEC)); + } +} \ No newline at end of file diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvAssertionsTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvAssertionsTest.java new file mode 100644 index 0000000000..b068738750 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvAssertionsTest.java @@ -0,0 +1,90 @@ +package com.dat3m.dartagnan.spirv.vulkan.ma; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.PROGRAM_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.FAIL; +import static com.dat3m.dartagnan.utils.Result.PASS; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvAssertionsTest { + + private final String modelPath = getRootPath("cat/vulkan.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvAssertionsTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/vulkan/ma/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"histogram-1.1.4.spvasm", 2, FAIL}, + {"histogram-2.1.2.spvasm", 2, PASS}, + {"histogram-4.1.1.spvasm", 2, PASS}, + {"histogram-lc2gb-1.spvasm", 2, FAIL}, + {"histogram-lc2gb-2.spvasm", 2, FAIL}, + {"compact-features-2.1.2.spvasm", 2, PASS}, + {"compact-features-lc2gb.spvasm", 2, FAIL}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.Z3); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.VULKAN); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(PROGRAM_SPEC)); + } +} diff --git a/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvRacesTest.java b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvRacesTest.java new file mode 100644 index 0000000000..0f723459e0 --- /dev/null +++ b/dartagnan/src/test/java/com/dat3m/dartagnan/spirv/vulkan/ma/SpirvRacesTest.java @@ -0,0 +1,91 @@ +package com.dat3m.dartagnan.spirv.vulkan.ma; + +import com.dat3m.dartagnan.configuration.Arch; +import com.dat3m.dartagnan.encoding.ProverWithTracker; +import com.dat3m.dartagnan.parsers.cat.ParserCat; +import com.dat3m.dartagnan.parsers.program.ProgramParser; +import com.dat3m.dartagnan.program.Program; +import com.dat3m.dartagnan.utils.Result; +import com.dat3m.dartagnan.verification.VerificationTask; +import com.dat3m.dartagnan.verification.solving.AssumeSolver; +import com.dat3m.dartagnan.wmm.Wmm; +import org.junit.Test; +import org.junit.runner.RunWith; +import org.junit.runners.Parameterized; +import org.sosy_lab.common.ShutdownManager; +import org.sosy_lab.common.configuration.Configuration; +import org.sosy_lab.common.configuration.InvalidConfigurationException; +import org.sosy_lab.common.log.BasicLogManager; +import org.sosy_lab.java_smt.SolverContextFactory; +import org.sosy_lab.java_smt.api.SolverContext; + +import java.io.File; +import java.io.IOException; +import java.util.Arrays; +import java.util.EnumSet; + +import static com.dat3m.dartagnan.configuration.Property.CAT_SPEC; +import static com.dat3m.dartagnan.utils.ResourceHelper.getRootPath; +import static com.dat3m.dartagnan.utils.ResourceHelper.getTestResourcePath; +import static com.dat3m.dartagnan.utils.Result.FAIL; +import static com.dat3m.dartagnan.utils.Result.PASS; +import static org.junit.Assert.assertEquals; + +@RunWith(Parameterized.class) +public class SpirvRacesTest { + + private final String modelPath = getRootPath("cat/vulkan.cat"); + private final String programPath; + private final int bound; + private final Result expected; + + public SpirvRacesTest(String file, int bound, Result expected) { + this.programPath = getTestResourcePath("spirv/vulkan/ma/" + file); + this.bound = bound; + this.expected = expected; + } + + @Parameterized.Parameters(name = "{index}: {0}, {1}, {2}") + public static Iterable data() throws IOException { + return Arrays.asList(new Object[][]{ + {"histogram-1.1.4.spvasm", 2, PASS}, + {"histogram-2.1.2.spvasm", 2, PASS}, + {"histogram-4.1.1.spvasm", 2, PASS}, + {"histogram-dv2wg.spvasm", 2, FAIL}, + {"histogram-lc2gb-1.spvasm", 2, FAIL}, + {"histogram-lc2gb-2.spvasm", 2, FAIL}, + {"compact-features-2.1.2.spvasm", 2, PASS}, + {"compact-features-lc2gb.spvasm", 2, FAIL}, + }); + } + + @Test + public void test() throws Exception { + try (SolverContext ctx = mkCtx(); ProverWithTracker prover = mkProver(ctx)) { + assertEquals(expected, AssumeSolver.run(ctx, prover, mkTask()).getResult()); + } + } + + private SolverContext mkCtx() throws InvalidConfigurationException { + Configuration cfg = Configuration.builder().build(); + return SolverContextFactory.createSolverContext( + cfg, + BasicLogManager.create(cfg), + ShutdownManager.create().getNotifier(), + SolverContextFactory.Solvers.YICES2); + } + + private ProverWithTracker mkProver(SolverContext ctx) { + return new ProverWithTracker(ctx, "", SolverContext.ProverOptions.GENERATE_MODELS); + } + + private VerificationTask mkTask() throws Exception { + VerificationTask.VerificationTaskBuilder builder = VerificationTask.builder() + .withConfig(Configuration.builder().build()) + .withBound(bound) + .withTarget(Arch.VULKAN); + Program program = new ProgramParser().parse(new File(programPath)); + Wmm mcm = new ParserCat().parse(new File(modelPath)); + return builder.build(program, mcm, EnumSet.of(CAT_SPEC)); + } +} \ No newline at end of file diff --git a/dartagnan/src/test/resources/OPENCL-DR-expected.csv b/dartagnan/src/test/resources/OPENCL-DR-expected.csv index 90b90623fc..9bddaf3c04 100644 --- a/dartagnan/src/test/resources/OPENCL-DR-expected.csv +++ b/dartagnan/src/test/resources/OPENCL-DR-expected.csv @@ -21,6 +21,10 @@ litmus/OPENCL/herd/old/MP_dr.litmus,0 litmus/OPENCL/herd/old/MP_relacq.litmus,0 litmus/OPENCL/herd/old/MP_relaxed.litmus,0 litmus/OPENCL/herd/old/MP_relseq.litmus,0 +litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/histgram.litmus,1 +litmus/OPENCL/mixedAtomicity/MP.litmus,0 +litmus/OPENCL/mixedAtomicity/partial-atomic.litmus,0 litmus/OPENCL/overhauling/example4.litmus,1 litmus/OPENCL/overhauling/example5.litmus,0 litmus/OPENCL/overhauling/example6.litmus,1 diff --git a/dartagnan/src/test/resources/OPENCL-MA-DR-expected.csv b/dartagnan/src/test/resources/OPENCL-MA-DR-expected.csv new file mode 100644 index 0000000000..ac27229260 --- /dev/null +++ b/dartagnan/src/test/resources/OPENCL-MA-DR-expected.csv @@ -0,0 +1,63 @@ +litmus/OPENCL/herd/2+2W.litmus,1 +litmus/OPENCL/herd/3.2W.litmus,0 +litmus/OPENCL/herd/3LB.litmus,1 +litmus/OPENCL/herd/CT_wsq1.litmus,1 +litmus/OPENCL/herd/IRIW.litmus,1 +litmus/OPENCL/herd/ISA2.litmus,0 +litmus/OPENCL/herd/ISA3.litmus,1 +litmus/OPENCL/herd/LB.litmus,0 +litmus/OPENCL/herd/MP.litmus,0 +litmus/OPENCL/herd/R.litmus,1 +litmus/OPENCL/herd/RWC.litmus,0 +litmus/OPENCL/herd/S.litmus,0 +litmus/OPENCL/herd/SB.litmus,0 +litmus/OPENCL/herd/SB1.litmus,0 +litmus/OPENCL/herd/SB2.litmus,0 +litmus/OPENCL/herd/WRC.litmus,0 +litmus/OPENCL/herd/barrier_example.litmus,1 +litmus/OPENCL/herd/global_barrier.litmus,0 +litmus/OPENCL/herd/global_barrier_mo.litmus,1 +litmus/OPENCL/herd/thinair.litmus,1 +litmus/OPENCL/herd/old/MP_dr.litmus,0 +litmus/OPENCL/herd/old/MP_relacq.litmus,0 +litmus/OPENCL/herd/old/MP_relaxed.litmus,0 +litmus/OPENCL/herd/old/MP_relseq.litmus,0 +litmus/OPENCL/overhauling/example4.litmus,1 +litmus/OPENCL/overhauling/example5.litmus,0 +litmus/OPENCL/overhauling/example6.litmus,1 +litmus/OPENCL/overhauling/example7a.litmus,1 +litmus/OPENCL/overhauling/example7b.litmus,0 +litmus/OPENCL/overhauling/example8.litmus,0 +litmus/OPENCL/overhauling/example9a.litmus,1 +litmus/OPENCL/overhauling/example9b.litmus,1 +litmus/OPENCL/overhauling/example10.litmus,1 +litmus/OPENCL/overhauling/IRIW_sc_dev.litmus,1 +litmus/OPENCL/overhauling/IRIW_sc_wg.litmus,1 +litmus/OPENCL/overhauling/ISA2.litmus,1 +litmus/OPENCL/overhauling/ISA2_broken.litmus,0 +litmus/OPENCL/overhauling/MP_ra_dev.litmus,1 +litmus/OPENCL/overhauling/MP_ra_dev_broken.litmus,0 +litmus/OPENCL/overhauling/MP_ra_wg.litmus,0 +litmus/OPENCL/overhauling/MP_ra_dev.litmus,1 +litmus/OPENCL/mixedAtomicity/histogram.litmus,1 +litmus/OPENCL/mixedAtomicity/rmw-bar.litmus,1 +litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/barrier-unordered.litmus,1 +litmus/OPENCL/mixedAtomicity/wkwk-mp-ordered.litmus,0 +litmus/OPENCL/mixedAtomicity/mp-an-relacq.litmus,1 +litmus/OPENCL/mixedAtomicity/mp-na-relacq.litmus,1 +litmus/OPENCL/mixedAtomicity/rlxrlx-mp-ordered.litmus,0 +litmus/OPENCL/mixedAtomicity/relrlx-mp-ordered.litmus,0 +litmus/OPENCL/mixedAtomicity/rlxacq-mp-ordered.litmus,0 +litmus/OPENCL/mixedAtomicity/fence-relacq.litmus,1 +litmus/OPENCL/mixedAtomicity/sc-fence-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/partial-atomic.litmus,0 +litmus/OPENCL/mixedAtomicity/partial-atomic1.litmus,0 +litmus/OPENCL/mixedAtomicity/partial-atomic2.litmus,1 +litmus/OPENCL/mixedAtomicity/mp-fence-rel-acq.litmus,1 +litmus/OPENCL/mixedAtomicity/ttas-lock.litmus,0 +litmus/OPENCL/mixedAtomicity/lb-nn.litmus,1 +litmus/OPENCL/mixedAtomicity/lb-an.litmus,1 +litmus/OPENCL/mixedAtomicity/co-ar.litmus,0 +litmus/OPENCL/mixedAtomicity/co-nr.litmus,0 +litmus/OPENCL/mixedAtomicity/corr.litmus,0 diff --git a/dartagnan/src/test/resources/OPENCL-MA-expected.csv b/dartagnan/src/test/resources/OPENCL-MA-expected.csv new file mode 100644 index 0000000000..53f592b676 --- /dev/null +++ b/dartagnan/src/test/resources/OPENCL-MA-expected.csv @@ -0,0 +1,189 @@ +litmus/OPENCL/herd/2+2W.litmus,0 +litmus/OPENCL/herd/3.2W.litmus,1 +litmus/OPENCL/herd/3LB.litmus,0 +litmus/OPENCL/herd/CT_wsq1.litmus,0 +litmus/OPENCL/herd/CT_wsq2.litmus,0 +litmus/OPENCL/herd/IRIW.litmus,0 +litmus/OPENCL/herd/ISA2.litmus,1 +litmus/OPENCL/herd/ISA3.litmus,0 +litmus/OPENCL/herd/LB.litmus,1 +litmus/OPENCL/herd/MP.litmus,1 +litmus/OPENCL/herd/R.litmus,1 +litmus/OPENCL/herd/RWC.litmus,1 +litmus/OPENCL/herd/S.litmus,1 +litmus/OPENCL/herd/SB.litmus,1 +litmus/OPENCL/herd/SB1.litmus,1 +litmus/OPENCL/herd/WRC.litmus,0 +litmus/OPENCL/herd/barrier_example.litmus,0 +litmus/OPENCL/herd/global_barrier.litmus,1 +litmus/OPENCL/herd/global_barrier_mo.litmus,1 +litmus/OPENCL/herd/thinair.litmus,1 +litmus/OPENCL/herd/old/MP_dr.litmus,1 +litmus/OPENCL/herd/old/MP_relacq.litmus,1 +litmus/OPENCL/herd/old/MP_relaxed.litmus,1 +litmus/OPENCL/herd/old/MP_relseq.litmus,1 +litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/barrier-unordered.litmus,1 +litmus/OPENCL/mixedAtomicity/histogram.litmus,1 +litmus/OPENCL/mixedAtomicity/rmw-bar.litmus,1 +litmus/OPENCL/mixedAtomicity/mp-an-relacq.litmus,1 +litmus/OPENCL/mixedAtomicity/sc-fence-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/partial-atomic2.litmus,1 +litmus/OPENCL/mixedAtomicity/mp-fence-rel-acq.litmus,1 +litmus/OPENCL/mixedAtomicity/ttas-lock.litmus,1 +litmus/OPENCL/mixedAtomicity/lb-nn.litmus,1 +litmus/OPENCL/mixedAtomicity/lb-an.litmus,1 +litmus/OPENCL/mixedAtomicity/fence-relacq.litmus,1 +litmus/OPENCL/overhauling/example4.litmus,0 +litmus/OPENCL/overhauling/example5.litmus,1 +litmus/OPENCL/overhauling/example6.litmus,0 +litmus/OPENCL/overhauling/example7a.litmus,0 +litmus/OPENCL/overhauling/example7b.litmus,1 +litmus/OPENCL/overhauling/example8.litmus,1 +litmus/OPENCL/overhauling/example9a.litmus,0 +litmus/OPENCL/overhauling/example9b.litmus,0 +litmus/OPENCL/overhauling/example10.litmus,1 +litmus/OPENCL/overhauling/IRIW_sc_dev.litmus,0 +litmus/OPENCL/overhauling/IRIW_sc_wg.litmus,0 +litmus/OPENCL/overhauling/ISA2.litmus,0 +litmus/OPENCL/overhauling/ISA2_broken.litmus,1 +litmus/OPENCL/overhauling/MP_ra_dev.litmus,0 +litmus/OPENCL/overhauling/MP_ra_dev_broken.litmus,1 +litmus/OPENCL/overhauling/MP_ra_wg.litmus,1 +litmus/OPENCL/overhauling/MP_ra_dev.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+sc+Racq.litmus,1 +litmus/OPENCL/portedFromC11/auto/linearisation.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+sc+Rrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/roachmotel2.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Rna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+sc+Rna.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+rlx+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wrel+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+rlx+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wna+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_pq_reorder.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rrlx+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/seq2.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+sc+Rsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Rsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+rlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+acq+Wsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rrlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+acq+Racq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+acq+Rrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rna+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Racq+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wsc+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+acq+Wna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wna+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rlx+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wrel+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rlx+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/arfna.litmus,0 +litmus/OPENCL/portedFromC11/auto/b_reorder+sc+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/rseq_weak.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Wrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+sc+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rsc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Wrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rel+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rel+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wrlx+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wrlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rna+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/rseq_weak2.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Wrel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Wrel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rsc+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3v2.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rel+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Wsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wsc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/cyc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Wna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a4_reorder.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Racq+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_q.litmus,0 +litmus/OPENCL/portedFromC11/auto/c.litmus,0 +litmus/OPENCL/portedFromC11/auto/a4.litmus,0 +litmus/OPENCL/portedFromC11/auto/b_reorder+sc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Racq+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wsc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_pq.litmus,0 +litmus/OPENCL/portedFromC11/auto/a1.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+acq+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wrel+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/roachmotel.litmus,0 +litmus/OPENCL/portedFromC11/auto/b+acq+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wrel+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/fig1.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Rrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Rna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Racq.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder+rlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Racq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rrlx+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/cyc_na.litmus,0 +litmus/OPENCL/portedFromC11/auto/strengthen2.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Rrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rsc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wna+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rrlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+rel+Rsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/strengthen.litmus,0 +litmus/OPENCL/portedFromC11/auto/linearisation2.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wsc+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_reorder.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+acq+Wrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wrlx+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+sc+rlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/arfna2.litmus,0 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Wsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+sc+Wsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/seq.litmus,0 +litmus/OPENCL/portedFromC11/auto/b+sc+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rsc+rel.litmus,1 +litmus/OPENCL/portedFromC11/auto/lb.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+sc+Wna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1_reorder+sc+Wna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Rna+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+acq+Wrel.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Rna+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_p_reorder.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+sc+Wrlx.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Racq+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+acq+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_q_reorder.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+acq+Rna.litmus,1 +litmus/OPENCL/portedFromC11/auto/a1+Wna+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/b.litmus,1 +litmus/OPENCL/portedFromC11/auto/c_p.litmus,0 +litmus/OPENCL/portedFromC11/auto/a3+acq+Rsc.litmus,1 +litmus/OPENCL/portedFromC11/auto/b_reorder.litmus,1 +litmus/OPENCL/portedFromC11/auto/b+sc+sc.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3_reorder+Wrlx+acq.litmus,1 +litmus/OPENCL/portedFromC11/auto/a3+sc+Wrel.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.2.litmus,0 +litmus/OPENCL/portedFromC11/manual/mp_fences.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-E3.10.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-E3.4.litmus,0 +litmus/OPENCL/portedFromC11/manual/RWC-sc-acq-sc-sc-sc.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-E3.8.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.8-alt.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.6.litmus,1 +litmus/OPENCL/portedFromC11/manual/mp_relacq.litmus,0 +litmus/OPENCL/portedFromC11/manual/iriw_sc.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-E3.3.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.1.litmus,0 +litmus/OPENCL/portedFromC11/manual/cppmem_iriw_relacq.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.5.litmus,1 +litmus/OPENCL/portedFromC11/manual/imm-E3.9.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-R2.litmus,0 +litmus/OPENCL/portedFromC11/manual/example1.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-R2-alt.litmus,1 +litmus/OPENCL/portedFromC11/manual/mp_relaxed.litmus,1 +litmus/OPENCL/portedFromC11/manual/IRIW-sc-sc-acq-sc-acq-sc.litmus,0 +litmus/OPENCL/portedFromC11/manual/imm-E3.7.litmus,0 diff --git a/dartagnan/src/test/resources/OPENCL-expected.csv b/dartagnan/src/test/resources/OPENCL-expected.csv index dbe3bce17a..1424cb2e6c 100644 --- a/dartagnan/src/test/resources/OPENCL-expected.csv +++ b/dartagnan/src/test/resources/OPENCL-expected.csv @@ -22,6 +22,9 @@ litmus/OPENCL/herd/old/MP_dr.litmus,1 litmus/OPENCL/herd/old/MP_relacq.litmus,1 litmus/OPENCL/herd/old/MP_relaxed.litmus,1 litmus/OPENCL/herd/old/MP_relseq.litmus,1 +litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus,1 +litmus/OPENCL/mixedAtomicity/histgram.litmus,1 +litmus/OPENCL/mixedAtomicity/MP.litmus,1 litmus/OPENCL/overhauling/example4.litmus,0 litmus/OPENCL/overhauling/example5.litmus,1 litmus/OPENCL/overhauling/example6.litmus,0 diff --git a/dartagnan/src/test/resources/PTXv7_5-expected.csv b/dartagnan/src/test/resources/PTXv7_5-expected.csv index 0fb5e6be99..fd3403ed59 100644 --- a/dartagnan/src/test/resources/PTXv7_5-expected.csv +++ b/dartagnan/src/test/resources/PTXv7_5-expected.csv @@ -184,6 +184,7 @@ litmus/PTX/Manual/SB+sc-gpu-multiFence-TotalOrder.litmus,1 litmus/PTX/Manual/SB+sc-sys.litmus,1 litmus/PTX/Manual/SB+sc-sys-gpu.litmus,1 litmus/PTX/Manual/SB-weak.litmus,1 +litmus/PTX/Manual/Mixed-Atomicity-barrier-ordered.litmus,1 litmus/PTX/Manual/Ticketlock-same-gpu.litmus,0 litmus/PTX/Manual/Ticketlock-diff-gpu.litmus,1 litmus/PTX/Manual/Ticketlock-acq2rlx-1.litmus,0 diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-2.1.2.spvasm new file mode 100644 index 0000000000..b338a7bda8 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-2.1.2.spvasm @@ -0,0 +1,156 @@ +; @Input: %flags_0 = {1, 1, 0, 1} +; @Input: %out_indices_0 = {42, 42, 42, 42} +; @Input: %group_offset_0 = {0, 2} +; @Output: forall ((%out_indices_0[0] == 0 and %out_indices_0[1] == 1) or (%out_indices_0[0] == 1 and %out_indices_0[1] == 0)) and (%out_indices_0[2] == 3 and %out_indices_0[3] == 42) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 66 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %60 "compact_features" %compact_features_s_idx %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInWorkgroupId + OpSource OpenCL_C 200000 + OpName %compact_features_s_idx "compact_features.s_idx" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %__spirv_BuiltInWorkgroupId "__spirv_BuiltInWorkgroupId" + OpName %compact_features "compact_features" + OpName %flags "flags" + OpName %out_indices "out_indices" + OpName %group_offset "group_offset" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %if_then4 "if.then4" + OpName %if_end7 "if.end7" + OpName %flags_addr "flags.addr" + OpName %out_indices_addr "out_indices.addr" + OpName %group_offset_addr "group_offset.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %group_id "group_id" + OpName %dst "dst" + OpName %call "call" + OpName %call1 "call1" + OpName %call2 "call2" + OpName %cmp "cmp" + OpName %arrayidx "arrayidx" + OpName %arrayidx3 "arrayidx3" + OpName %tobool "tobool" + OpName %call5 "call5" + OpName %arrayidx6 "arrayidx6" + OpName %flags_0 "flags" + OpName %out_indices_0 "out_indices" + OpName %group_offset_0 "group_offset" + OpDecorate %compact_features_s_idx Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInWorkgroupId LinkageAttributes "__spirv_BuiltInWorkgroupId" Import + OpDecorate %__spirv_BuiltInWorkgroupId Constant + OpDecorate %__spirv_BuiltInWorkgroupId BuiltIn WorkgroupId + OpDecorate %compact_features LinkageAttributes "compact_features" Export + OpDecorate %flags Alignment 4 + OpDecorate %out_indices Alignment 4 + OpDecorate %group_offset Alignment 4 + OpDecorate %flags_addr Alignment 4 + OpDecorate %out_indices_addr Alignment 4 + OpDecorate %group_offset_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %group_id Alignment 4 + OpDecorate %dst Alignment 4 + OpDecorate %flags_0 Alignment 4 + OpDecorate %out_indices_0 Alignment 4 + OpDecorate %group_offset_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %12 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint + %bool = OpTypeBool +%_ptr_Generic_uint = OpTypePointer Generic %uint +%compact_features_s_idx = OpVariable %_ptr_Workgroup_uint Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInWorkgroupId = OpVariable %_ptr_Input_v3uint Input +%compact_features = OpFunction %void DontInline %12 + %flags = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %flags_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%out_indices_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%group_offset_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %group_id = OpVariable %_ptr_Function_uint Function + %dst = OpVariable %_ptr_Function_uint Function + OpStore %flags_addr %flags Aligned 4 + OpStore %out_indices_addr %out_indices Aligned 4 + OpStore %group_offset_addr %group_offset Aligned 4 + %31 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %31 0 + OpStore %tid %call Aligned 4 + %33 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %33 0 + OpStore %gid %call1 Aligned 4 + %35 = OpLoad %v3uint %__spirv_BuiltInWorkgroupId Aligned 16 + %call2 = OpCompositeExtract %uint %35 0 + OpStore %group_id %call2 Aligned 4 + %37 = OpLoad %uint %tid Aligned 4 + %cmp = OpIEqual %bool %37 %uint_0 + %42 = OpPtrCastToGeneric %_ptr_Generic_uint %compact_features_s_idx + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %43 = OpLoad %_ptr_CrossWorkgroup_uint %group_offset_addr Aligned 4 + %44 = OpLoad %uint %group_id Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %43 %44 + %46 = OpLoad %uint %arrayidx Aligned 4 + OpStore %compact_features_s_idx %46 Aligned 4 + OpBranch %if_end + %if_end = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_272 + %49 = OpLoad %_ptr_CrossWorkgroup_uint %flags_addr Aligned 4 + %50 = OpLoad %uint %gid Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %49 %50 + %52 = OpLoad %uint %arrayidx3 Aligned 4 + %tobool = OpINotEqual %bool %52 %uint_0 + OpBranchConditional %tobool %if_then4 %if_end7 + %if_then4 = OpLabel + %call5 = OpAtomicIAdd %uint %42 %uint_1 %uint_0 %uint_1 + OpStore %dst %call5 Aligned 4 + %56 = OpLoad %uint %gid Aligned 4 + %57 = OpLoad %_ptr_CrossWorkgroup_uint %out_indices_addr Aligned 4 + %58 = OpLoad %uint %dst Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %57 %58 + OpStore %arrayidx6 %56 Aligned 4 + OpBranch %if_end7 + %if_end7 = OpLabel + OpReturn + OpFunctionEnd + %60 = OpFunction %void DontInline %12 + %flags_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %64 = OpLabel + %65 = OpFunctionCall %void %compact_features %flags_0 %out_indices_0 %group_offset_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-implicit-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-implicit-2.1.2.spvasm new file mode 100644 index 0000000000..f25f613af7 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-implicit-2.1.2.spvasm @@ -0,0 +1,152 @@ +; @Input: %flags_0 = {1, 1, 0, 1} +; @Input: %out_indices_0 = {42, 42, 42, 42} +; @Input: %group_offset_0 = {0, 2} +; @Output: forall ((%out_indices_0[0] == 0 and %out_indices_0[1] == 1) or (%out_indices_0[0] == 1 and %out_indices_0[1] == 0)) and (%out_indices_0[2] == 3 and %out_indices_0[3] == 42) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 63 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %57 "compact_features" %compact_features_s_idx %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInWorkgroupId + OpSource OpenCL_C 200000 + OpName %compact_features_s_idx "compact_features.s_idx" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %__spirv_BuiltInWorkgroupId "__spirv_BuiltInWorkgroupId" + OpName %compact_features "compact_features" + OpName %flags "flags" + OpName %out_indices "out_indices" + OpName %group_offset "group_offset" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %if_then4 "if.then4" + OpName %if_end7 "if.end7" + OpName %flags_addr "flags.addr" + OpName %out_indices_addr "out_indices.addr" + OpName %group_offset_addr "group_offset.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %group_id "group_id" + OpName %dst "dst" + OpName %call "call" + OpName %call1 "call1" + OpName %call2 "call2" + OpName %cmp "cmp" + OpName %arrayidx "arrayidx" + OpName %arrayidx3 "arrayidx3" + OpName %tobool "tobool" + OpName %call5 "call5" + OpName %arrayidx6 "arrayidx6" + OpName %flags_0 "flags" + OpName %out_indices_0 "out_indices" + OpName %group_offset_0 "group_offset" + OpDecorate %compact_features_s_idx Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInWorkgroupId LinkageAttributes "__spirv_BuiltInWorkgroupId" Import + OpDecorate %__spirv_BuiltInWorkgroupId Constant + OpDecorate %__spirv_BuiltInWorkgroupId BuiltIn WorkgroupId + OpDecorate %compact_features LinkageAttributes "compact_features" Export + OpDecorate %flags Alignment 4 + OpDecorate %out_indices Alignment 4 + OpDecorate %group_offset Alignment 4 + OpDecorate %flags_addr Alignment 4 + OpDecorate %out_indices_addr Alignment 4 + OpDecorate %group_offset_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %group_id Alignment 4 + OpDecorate %dst Alignment 4 + OpDecorate %flags_0 Alignment 4 + OpDecorate %out_indices_0 Alignment 4 + OpDecorate %group_offset_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %uint_272 = OpConstant %uint 272 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %12 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint + %bool = OpTypeBool +%compact_features_s_idx = OpVariable %_ptr_Workgroup_uint Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInWorkgroupId = OpVariable %_ptr_Input_v3uint Input +%compact_features = OpFunction %void DontInline %12 + %flags = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %flags_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%out_indices_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%group_offset_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %group_id = OpVariable %_ptr_Function_uint Function + %dst = OpVariable %_ptr_Function_uint Function + OpStore %flags_addr %flags Aligned 4 + OpStore %out_indices_addr %out_indices Aligned 4 + OpStore %group_offset_addr %group_offset Aligned 4 + %31 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %31 0 + OpStore %tid %call Aligned 4 + %33 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %33 0 + OpStore %gid %call1 Aligned 4 + %35 = OpLoad %v3uint %__spirv_BuiltInWorkgroupId Aligned 16 + %call2 = OpCompositeExtract %uint %35 0 + OpStore %group_id %call2 Aligned 4 + %37 = OpLoad %uint %tid Aligned 4 + %cmp = OpIEqual %bool %37 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %41 = OpLoad %_ptr_CrossWorkgroup_uint %group_offset_addr Aligned 4 + %42 = OpLoad %uint %group_id Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %41 %42 + %44 = OpLoad %uint %arrayidx Aligned 4 + OpStore %compact_features_s_idx %44 Aligned 4 + OpBranch %if_end + %if_end = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_272 + %47 = OpLoad %_ptr_CrossWorkgroup_uint %flags_addr Aligned 4 + %48 = OpLoad %uint %gid Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %47 %48 + %50 = OpLoad %uint %arrayidx3 Aligned 4 + %tobool = OpINotEqual %bool %50 %uint_0 + OpBranchConditional %tobool %if_then4 %if_end7 + %if_then4 = OpLabel + %call5 = OpAtomicIIncrement %uint %compact_features_s_idx %uint_2 %uint_0 + OpStore %dst %call5 Aligned 4 + %53 = OpLoad %uint %gid Aligned 4 + %54 = OpLoad %_ptr_CrossWorkgroup_uint %out_indices_addr Aligned 4 + %55 = OpLoad %uint %dst Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %54 %55 + OpStore %arrayidx6 %53 Aligned 4 + OpBranch %if_end7 + %if_end7 = OpLabel + OpReturn + OpFunctionEnd + %57 = OpFunction %void DontInline %12 + %flags_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %61 = OpLabel + %62 = OpFunctionCall %void %compact_features %flags_0 %out_indices_0 %group_offset_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-lc2gb.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-lc2gb.spvasm new file mode 100644 index 0000000000..4f94155034 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/compact-features-lc2gb.spvasm @@ -0,0 +1,156 @@ +; @Input: %flags_0 = {1, 1, 0, 1} +; @Input: %out_indices_0 = {42, 42, 42, 42} +; @Input: %group_offset_0 = {0, 2} +; @Output: forall ((%out_indices_0[0] == 0 and %out_indices_0[1] == 1) or (%out_indices_0[0] == 1 and %out_indices_0[1] == 0)) and (%out_indices_0[2] == 3 and %out_indices_0[3] == 42) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 66 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %60 "compact_features" %compact_features_s_idx %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId %__spirv_BuiltInWorkgroupId + OpSource OpenCL_C 200000 + OpName %compact_features_s_idx "compact_features.s_idx" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %__spirv_BuiltInWorkgroupId "__spirv_BuiltInWorkgroupId" + OpName %compact_features "compact_features" + OpName %flags "flags" + OpName %out_indices "out_indices" + OpName %group_offset "group_offset" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %if_then4 "if.then4" + OpName %if_end7 "if.end7" + OpName %flags_addr "flags.addr" + OpName %out_indices_addr "out_indices.addr" + OpName %group_offset_addr "group_offset.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %group_id "group_id" + OpName %dst "dst" + OpName %call "call" + OpName %call1 "call1" + OpName %call2 "call2" + OpName %cmp "cmp" + OpName %arrayidx "arrayidx" + OpName %arrayidx3 "arrayidx3" + OpName %tobool "tobool" + OpName %call5 "call5" + OpName %arrayidx6 "arrayidx6" + OpName %flags_0 "flags" + OpName %out_indices_0 "out_indices" + OpName %group_offset_0 "group_offset" + OpDecorate %compact_features_s_idx Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %__spirv_BuiltInWorkgroupId LinkageAttributes "__spirv_BuiltInWorkgroupId" Import + OpDecorate %__spirv_BuiltInWorkgroupId Constant + OpDecorate %__spirv_BuiltInWorkgroupId BuiltIn WorkgroupId + OpDecorate %compact_features LinkageAttributes "compact_features" Export + OpDecorate %flags Alignment 4 + OpDecorate %out_indices Alignment 4 + OpDecorate %group_offset Alignment 4 + OpDecorate %flags_addr Alignment 4 + OpDecorate %out_indices_addr Alignment 4 + OpDecorate %group_offset_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %group_id Alignment 4 + OpDecorate %dst Alignment 4 + OpDecorate %flags_0 Alignment 4 + OpDecorate %out_indices_0 Alignment 4 + OpDecorate %group_offset_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_0 = OpConstant %uint 0 + %uint_2 = OpConstant %uint 2 + %uint_528 = OpConstant %uint 528 + %uint_1 = OpConstant %uint 1 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %12 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint + %bool = OpTypeBool +%_ptr_Generic_uint = OpTypePointer Generic %uint +%compact_features_s_idx = OpVariable %_ptr_Workgroup_uint Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInWorkgroupId = OpVariable %_ptr_Input_v3uint Input +%compact_features = OpFunction %void DontInline %12 + %flags = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel + %flags_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%out_indices_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%group_offset_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %group_id = OpVariable %_ptr_Function_uint Function + %dst = OpVariable %_ptr_Function_uint Function + OpStore %flags_addr %flags Aligned 4 + OpStore %out_indices_addr %out_indices Aligned 4 + OpStore %group_offset_addr %group_offset Aligned 4 + %31 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %31 0 + OpStore %tid %call Aligned 4 + %33 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %33 0 + OpStore %gid %call1 Aligned 4 + %35 = OpLoad %v3uint %__spirv_BuiltInWorkgroupId Aligned 16 + %call2 = OpCompositeExtract %uint %35 0 + OpStore %group_id %call2 Aligned 4 + %37 = OpLoad %uint %tid Aligned 4 + %cmp = OpIEqual %bool %37 %uint_0 + %42 = OpPtrCastToGeneric %_ptr_Generic_uint %compact_features_s_idx + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %43 = OpLoad %_ptr_CrossWorkgroup_uint %group_offset_addr Aligned 4 + %44 = OpLoad %uint %group_id Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %43 %44 + %46 = OpLoad %uint %arrayidx Aligned 4 + OpStore %compact_features_s_idx %46 Aligned 4 + OpBranch %if_end + %if_end = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_528 + %49 = OpLoad %_ptr_CrossWorkgroup_uint %flags_addr Aligned 4 + %50 = OpLoad %uint %gid Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %49 %50 + %52 = OpLoad %uint %arrayidx3 Aligned 4 + %tobool = OpINotEqual %bool %52 %uint_0 + OpBranchConditional %tobool %if_then4 %if_end7 + %if_then4 = OpLabel + %call5 = OpAtomicIAdd %uint %42 %uint_1 %uint_0 %uint_1 + OpStore %dst %call5 Aligned 4 + %56 = OpLoad %uint %gid Aligned 4 + %57 = OpLoad %_ptr_CrossWorkgroup_uint %out_indices_addr Aligned 4 + %58 = OpLoad %uint %dst Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %57 %58 + OpStore %arrayidx6 %56 Aligned 4 + OpBranch %if_end7 + %if_end7 = OpLabel + OpReturn + OpFunctionEnd + %60 = OpFunction %void DontInline %12 + %flags_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%out_indices_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%group_offset_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %64 = OpLabel + %65 = OpFunctionCall %void %compact_features %flags_0 %out_indices_0 %group_offset_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-1.1.4.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-1.1.4.spvasm new file mode 100644 index 0000000000..4601d8b07b --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-1.1.4.spvasm @@ -0,0 +1,141 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 1, 1, 4 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 64 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %59 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %47 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %47 + %49 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %49 Aligned 4 + %50 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %50 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %53 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %54 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %53 %54 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %57 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_1 %uint_0 %57 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %59 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %62 = OpLabel + %63 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-2.1.2.spvasm new file mode 100644 index 0000000000..43b4334ee7 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-2.1.2.spvasm @@ -0,0 +1,141 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 64 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %59 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %47 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %47 + %49 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %49 Aligned 4 + %50 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %50 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %53 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %54 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %53 %54 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %57 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_1 %uint_0 %57 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %59 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %62 = OpLabel + %63 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-4.1.1.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-4.1.1.spvasm new file mode 100644 index 0000000000..859772664b --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-4.1.1.spvasm @@ -0,0 +1,141 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 4, 1, 1 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 64 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %59 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %47 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %47 + %49 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %49 Aligned 4 + %50 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %50 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %53 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %54 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %53 %54 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %57 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_1 %uint_0 %57 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %59 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %62 = OpLabel + %63 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-dv2wg.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-dv2wg.spvasm new file mode 100644 index 0000000000..3325fdb67c --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-dv2wg.spvasm @@ -0,0 +1,141 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 64 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %59 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %47 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %47 + %49 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %49 Aligned 4 + %50 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %50 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %53 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %54 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %53 %54 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %57 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_2 %uint_0 %57 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %59 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %62 = OpLabel + %63 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-1.1.4.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-1.1.4.spvasm new file mode 100644 index 0000000000..c174bccd63 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-1.1.4.spvasm @@ -0,0 +1,136 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 1, 1, 4 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 62 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %57 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %add_ptr "add.ptr" + OpName %call3 "call3" + OpName %arrayidx4 "arrayidx4" + OpName %cmp "cmp" + OpName %add_ptr5 "add.ptr5" + OpName %call6 "call6" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %42 = OpBitcast %_ptr_Workgroup_uint %histo_main_kernel_sub_histo + %add_ptr = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %42 %41 + %call3 = OpAtomicIAdd %uint %add_ptr %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %46 = OpLoad %uint %tid Aligned 4 + %arrayidx4 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %46 + %48 = OpLoad %uint %arrayidx4 Aligned 4 + OpStore %count %48 Aligned 4 + %49 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %49 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %52 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %53 = OpLoad %uint %tid Aligned 4 + %add_ptr5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %52 %53 + %55 = OpLoad %uint %count Aligned 4 + %call6 = OpAtomicIAdd %uint %add_ptr5 %uint_2 %uint_0 %55 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %57 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %60 = OpLabel + %61 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-2.1.2.spvasm new file mode 100644 index 0000000000..3b908d67ec --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-2.1.2.spvasm @@ -0,0 +1,136 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 62 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %57 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %add_ptr "add.ptr" + OpName %call3 "call3" + OpName %arrayidx4 "arrayidx4" + OpName %cmp "cmp" + OpName %add_ptr5 "add.ptr5" + OpName %call6 "call6" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %42 = OpBitcast %_ptr_Workgroup_uint %histo_main_kernel_sub_histo + %add_ptr = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %42 %41 + %call3 = OpAtomicIAdd %uint %add_ptr %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %46 = OpLoad %uint %tid Aligned 4 + %arrayidx4 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %46 + %48 = OpLoad %uint %arrayidx4 Aligned 4 + OpStore %count %48 Aligned 4 + %49 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %49 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %52 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %53 = OpLoad %uint %tid Aligned 4 + %add_ptr5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %52 %53 + %55 = OpLoad %uint %count Aligned 4 + %call6 = OpAtomicIAdd %uint %add_ptr5 %uint_2 %uint_0 %55 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %57 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %60 = OpLabel + %61 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-4.1.1.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-4.1.1.spvasm new file mode 100644 index 0000000000..6771dff58e --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-implicit-4.1.1.spvasm @@ -0,0 +1,136 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 4, 1, 1 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 62 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %57 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %add_ptr "add.ptr" + OpName %call3 "call3" + OpName %arrayidx4 "arrayidx4" + OpName %cmp "cmp" + OpName %add_ptr5 "add.ptr5" + OpName %call6 "call6" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %42 = OpBitcast %_ptr_Workgroup_uint %histo_main_kernel_sub_histo + %add_ptr = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %42 %41 + %call3 = OpAtomicIAdd %uint %add_ptr %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %46 = OpLoad %uint %tid Aligned 4 + %arrayidx4 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %46 + %48 = OpLoad %uint %arrayidx4 Aligned 4 + OpStore %count %48 Aligned 4 + %49 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %49 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %52 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %53 = OpLoad %uint %tid Aligned 4 + %add_ptr5 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %52 %53 + %55 = OpLoad %uint %count Aligned 4 + %call6 = OpAtomicIAdd %uint %add_ptr5 %uint_2 %uint_0 %55 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %57 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %60 = OpLabel + %61 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-1.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-1.spvasm new file mode 100644 index 0000000000..012a3d79bd --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-1.spvasm @@ -0,0 +1,142 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 65 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %60 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_528 = OpConstant %uint 528 + %uint_1 = OpConstant %uint 1 + %uint_272 = OpConstant %uint 272 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_528 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %48 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %48 + %50 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %50 Aligned 4 + %51 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %51 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %54 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %55 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %54 %55 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %58 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_1 %uint_0 %58 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %60 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %63 = OpLabel + %64 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-2.spvasm b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-2.spvasm new file mode 100644 index 0000000000..90be74e9bf --- /dev/null +++ b/dartagnan/src/test/resources/spirv/opencl/ma/histogram-lc2gb-2.spvasm @@ -0,0 +1,142 @@ +; @Input: %sm_mappings_0 = {0, 1, 1, 1} +; @Input: %global_histo_0 = {0, 0} +; @Output: forall (%global_histo_0[0] == 1 and %global_histo_0[1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.4 +; Generator: Khronos LLVM/SPIR-V Translator; 14 +; Bound: 65 +; Schema: 0 + OpCapability Addresses + OpCapability Linkage + OpCapability Kernel + OpCapability GenericPointer + %1 = OpExtInstImport "OpenCL.std" + OpMemoryModel Physical32 OpenCL + OpEntryPoint Kernel %60 "histo_main_kernel" %histo_main_kernel_sub_histo %__spirv_BuiltInLocalInvocationId %__spirv_BuiltInGlobalInvocationId + OpSource OpenCL_C 200000 + OpName %histo_main_kernel_sub_histo "histo_main_kernel.sub_histo" + OpName %__spirv_BuiltInLocalInvocationId "__spirv_BuiltInLocalInvocationId" + OpName %__spirv_BuiltInGlobalInvocationId "__spirv_BuiltInGlobalInvocationId" + OpName %histo_main_kernel "histo_main_kernel" + OpName %sm_mappings "sm_mappings" + OpName %global_histo "global_histo" + OpName %entry "entry" + OpName %if_then "if.then" + OpName %if_end "if.end" + OpName %sm_mappings_addr "sm_mappings.addr" + OpName %global_histo_addr "global_histo.addr" + OpName %tid "tid" + OpName %gid "gid" + OpName %bin_index "bin_index" + OpName %count "count" + OpName %call "call" + OpName %call1 "call1" + OpName %arrayidx "arrayidx" + OpName %arrayidx2 "arrayidx2" + OpName %arrayidx3 "arrayidx3" + OpName %arrayidx3_ascast "arrayidx3.ascast" + OpName %call4 "call4" + OpName %arrayidx5 "arrayidx5" + OpName %cmp "cmp" + OpName %arrayidx6 "arrayidx6" + OpName %arrayidx6_ascast "arrayidx6.ascast" + OpName %call7 "call7" + OpName %sm_mappings_0 "sm_mappings" + OpName %global_histo_0 "global_histo" + OpDecorate %histo_main_kernel_sub_histo Alignment 4 + OpDecorate %__spirv_BuiltInLocalInvocationId LinkageAttributes "__spirv_BuiltInLocalInvocationId" Import + OpDecorate %__spirv_BuiltInLocalInvocationId Constant + OpDecorate %__spirv_BuiltInLocalInvocationId BuiltIn LocalInvocationId + OpDecorate %__spirv_BuiltInGlobalInvocationId LinkageAttributes "__spirv_BuiltInGlobalInvocationId" Import + OpDecorate %__spirv_BuiltInGlobalInvocationId Constant + OpDecorate %__spirv_BuiltInGlobalInvocationId BuiltIn GlobalInvocationId + OpDecorate %histo_main_kernel LinkageAttributes "histo_main_kernel" Export + OpDecorate %sm_mappings Alignment 4 + OpDecorate %global_histo Alignment 4 + OpDecorate %sm_mappings_addr Alignment 4 + OpDecorate %global_histo_addr Alignment 4 + OpDecorate %tid Alignment 4 + OpDecorate %gid Alignment 4 + OpDecorate %bin_index Alignment 4 + OpDecorate %count Alignment 4 + OpDecorate %sm_mappings_0 Alignment 4 + OpDecorate %global_histo_0 Alignment 4 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %uint_0 = OpConstant %uint 0 + %uint_272 = OpConstant %uint 272 + %uint_1 = OpConstant %uint 1 + %uint_528 = OpConstant %uint 528 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %void = OpTypeVoid +%_ptr_CrossWorkgroup_uint = OpTypePointer CrossWorkgroup %uint + %13 = OpTypeFunction %void %_ptr_CrossWorkgroup_uint %_ptr_CrossWorkgroup_uint +%_ptr_Function__ptr_CrossWorkgroup_uint = OpTypePointer Function %_ptr_CrossWorkgroup_uint +%_ptr_Function_uint = OpTypePointer Function %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint +%_ptr_Generic_uint = OpTypePointer Generic %uint + %bool = OpTypeBool +%histo_main_kernel_sub_histo = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup +%__spirv_BuiltInLocalInvocationId = OpVariable %_ptr_Input_v3uint Input +%__spirv_BuiltInGlobalInvocationId = OpVariable %_ptr_Input_v3uint Input +%histo_main_kernel = OpFunction %void DontInline %13 +%sm_mappings = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %entry = OpLabel +%sm_mappings_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function +%global_histo_addr = OpVariable %_ptr_Function__ptr_CrossWorkgroup_uint Function + %tid = OpVariable %_ptr_Function_uint Function + %gid = OpVariable %_ptr_Function_uint Function + %bin_index = OpVariable %_ptr_Function_uint Function + %count = OpVariable %_ptr_Function_uint Function + OpStore %sm_mappings_addr %sm_mappings Aligned 4 + OpStore %global_histo_addr %global_histo Aligned 4 + %28 = OpLoad %v3uint %__spirv_BuiltInLocalInvocationId Aligned 16 + %call = OpCompositeExtract %uint %28 0 + OpStore %tid %call Aligned 4 + %30 = OpLoad %v3uint %__spirv_BuiltInGlobalInvocationId Aligned 16 + %call1 = OpCompositeExtract %uint %30 0 + OpStore %gid %call1 Aligned 4 + %32 = OpLoad %uint %tid Aligned 4 + %arrayidx = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %32 + OpStore %arrayidx %uint_0 Aligned 4 + OpControlBarrier %uint_2 %uint_2 %uint_272 + %37 = OpLoad %_ptr_CrossWorkgroup_uint %sm_mappings_addr Aligned 4 + %38 = OpLoad %uint %gid Aligned 4 + %arrayidx2 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %37 %38 + %40 = OpLoad %uint %arrayidx2 Aligned 4 + OpStore %bin_index %40 Aligned 4 + %41 = OpLoad %uint %bin_index Aligned 4 + %arrayidx3 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %41 +%arrayidx3_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx3 + %call4 = OpAtomicIAdd %uint %arrayidx3_ascast %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_528 + %48 = OpLoad %uint %tid Aligned 4 + %arrayidx5 = OpInBoundsPtrAccessChain %_ptr_Workgroup_uint %histo_main_kernel_sub_histo %uint_0 %48 + %50 = OpLoad %uint %arrayidx5 Aligned 4 + OpStore %count %50 Aligned 4 + %51 = OpLoad %uint %count Aligned 4 + %cmp = OpUGreaterThan %bool %51 %uint_0 + OpBranchConditional %cmp %if_then %if_end + %if_then = OpLabel + %54 = OpLoad %_ptr_CrossWorkgroup_uint %global_histo_addr Aligned 4 + %55 = OpLoad %uint %tid Aligned 4 + %arrayidx6 = OpInBoundsPtrAccessChain %_ptr_CrossWorkgroup_uint %54 %55 +%arrayidx6_ascast = OpPtrCastToGeneric %_ptr_Generic_uint %arrayidx6 + %58 = OpLoad %uint %count Aligned 4 + %call7 = OpAtomicIAdd %uint %arrayidx6_ascast %uint_1 %uint_0 %58 + OpBranch %if_end + %if_end = OpLabel + OpReturn + OpFunctionEnd + %60 = OpFunction %void DontInline %13 +%sm_mappings_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint +%global_histo_0 = OpFunctionParameter %_ptr_CrossWorkgroup_uint + %63 = OpLabel + %64 = OpFunctionCall %void %histo_main_kernel %sm_mappings_0 %global_histo_0 + OpReturn + OpFunctionEnd diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-2.1.2.spvasm new file mode 100644 index 0000000000..ffcb2f3bcd --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-2.1.2.spvasm @@ -0,0 +1,127 @@ +; @Input: %21 = {{1, 1, 0, 1}} +; @Input: %22 = {{42, 42, 42, 42}} +; @Input: %23 = {{0, 2}} +; @Output: forall ((%22[0][0] == 0 and %22[0][1] == 1) or (%22[0][0] == 1 and %22[0][1] == 0)) and (%22[0][2] == 3 and %22[0][3] == 42) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 88 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %66 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "compact_features" %3 %gl_GlobalInvocationID %gl_LocalInvocationID %gl_WorkGroupID %17 %21 %22 %23 %7 + OpSource OpenCL_C 200 + %67 = OpString "compact_features" + %68 = OpString "__kernel" + %71 = OpString "flags" + %74 = OpString "out_indices" + %77 = OpString "group_offset" + OpMemberDecorate %_struct_5 0 Offset 0 + OpMemberDecorate %_struct_5 1 Offset 16 + OpDecorate %_struct_5 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_19 0 Offset 0 + OpDecorate %_struct_19 Block + OpDecorate %21 DescriptorSet 0 + OpDecorate %21 Binding 0 + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 1 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 2 + OpDecorate %12 SpecId 0 + OpDecorate %13 SpecId 1 + OpDecorate %14 SpecId 2 + %uint = OpTypeInt 32 0 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %v3uint = OpTypeVector %uint 3 + %_struct_5 = OpTypeStruct %v3uint %v3uint +%_ptr_PushConstant__struct_5 = OpTypePointer PushConstant %_struct_5 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %12 = OpSpecConstant %uint 1 + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %12 %13 %14 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_19 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_19 = OpTypePointer StorageBuffer %_struct_19 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_2 = OpConstant %uint 2 + %uint_264 = OpConstant %uint 264 + %uint_3 = OpConstant %uint 3 + %uint_12 = OpConstant %uint 12 + %uint_16 = OpConstant %uint 16 + %uint_4 = OpConstant %uint 4 + %3 = OpVariable %_ptr_Workgroup_uint Workgroup + %7 = OpVariable %_ptr_PushConstant__struct_5 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input + %17 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %21 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %22 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %7 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %37 = OpIAdd %uint %36 %33 + %39 = OpAccessChain %_ptr_PushConstant_uint %7 %uint_1 %uint_0 + %40 = OpLoad %uint %39 Aligned 16 + %42 = OpIEqual %bool %31 %uint_0 + OpSelectionMerge %53 None + OpBranchConditional %42 %45 %53 + %45 = OpLabel + %46 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0 + %47 = OpLoad %uint %46 Aligned 16 + %48 = OpIAdd %uint %47 %40 + %50 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %48 + %51 = OpLoad %uint %50 Aligned 4 + OpStore %3 %51 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpBranch %53 + %53 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_264 + %56 = OpAccessChain %_ptr_StorageBuffer_uint %21 %uint_0 %37 + %57 = OpLoad %uint %56 Aligned 4 + %58 = OpINotEqual %bool %57 %uint_0 + OpSelectionMerge %65 None + OpBranchConditional %58 %61 %65 + %61 = OpLabel + %62 = OpAtomicIAdd %uint %3 %uint_5 %uint_0 %uint_1 + %63 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %62 + OpStore %63 %37 Aligned 4 + OpBranch %65 + %65 = OpLabel + OpReturn + OpFunctionEnd + %81 = OpExtInst %void %66 PushConstantRegionOffset %uint_0 %uint_12 + %83 = OpExtInst %void %66 PushConstantRegionGroupOffset %uint_16 %uint_12 + %70 = OpExtInst %void %66 Kernel %26 %67 %uint_3 %uint_0 %68 + %72 = OpExtInst %void %66 ArgumentInfo %71 + %73 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_0 %uint_0 %uint_0 %72 + %75 = OpExtInst %void %66 ArgumentInfo %74 + %76 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_1 %uint_0 %uint_1 %75 + %78 = OpExtInst %void %66 ArgumentInfo %77 + %79 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_2 %uint_0 %uint_2 %78 + %84 = OpExtInst %void %66 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %86 = OpExtInst %void %66 WorkgroupVariableSize %3 %uint_4 \ No newline at end of file diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-lc2gb.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-lc2gb.spvasm new file mode 100644 index 0000000000..5eef2fdb4c --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/compact-features-lc2gb.spvasm @@ -0,0 +1,127 @@ +; @Input: %21 = {{1, 1, 0, 1}} +; @Input: %22 = {{42, 42, 42, 42}} +; @Input: %23 = {{0, 2}} +; @Output: forall ((%22[0][0] == 0 and %22[0][1] == 1) or (%22[0][0] == 1 and %22[0][1] == 0)) and (%22[0][2] == 3 and %22[0][3] == 42) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 88 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %66 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "compact_features" %3 %gl_GlobalInvocationID %gl_LocalInvocationID %gl_WorkGroupID %17 %21 %22 %23 %7 + OpSource OpenCL_C 200 + %67 = OpString "compact_features" + %68 = OpString "__kernel" + %71 = OpString "flags" + %74 = OpString "out_indices" + %77 = OpString "group_offset" + OpMemberDecorate %_struct_5 0 Offset 0 + OpMemberDecorate %_struct_5 1 Offset 16 + OpDecorate %_struct_5 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_19 0 Offset 0 + OpDecorate %_struct_19 Block + OpDecorate %21 DescriptorSet 0 + OpDecorate %21 Binding 0 + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 1 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 2 + OpDecorate %12 SpecId 0 + OpDecorate %13 SpecId 1 + OpDecorate %14 SpecId 2 + %uint = OpTypeInt 32 0 +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %v3uint = OpTypeVector %uint 3 + %_struct_5 = OpTypeStruct %v3uint %v3uint +%_ptr_PushConstant__struct_5 = OpTypePointer PushConstant %_struct_5 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %12 = OpSpecConstant %uint 1 + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %12 %13 %14 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_19 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_19 = OpTypePointer StorageBuffer %_struct_19 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_2 = OpConstant %uint 2 + %uint_72 = OpConstant %uint 72 + %uint_3 = OpConstant %uint 3 + %uint_12 = OpConstant %uint 12 + %uint_16 = OpConstant %uint 16 + %uint_4 = OpConstant %uint 4 + %3 = OpVariable %_ptr_Workgroup_uint Workgroup + %7 = OpVariable %_ptr_PushConstant__struct_5 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input + %17 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %21 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %22 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_19 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %7 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %37 = OpIAdd %uint %36 %33 + %39 = OpAccessChain %_ptr_PushConstant_uint %7 %uint_1 %uint_0 + %40 = OpLoad %uint %39 Aligned 16 + %42 = OpIEqual %bool %31 %uint_0 + OpSelectionMerge %53 None + OpBranchConditional %42 %45 %53 + %45 = OpLabel + %46 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0 + %47 = OpLoad %uint %46 Aligned 16 + %48 = OpIAdd %uint %47 %40 + %50 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %48 + %51 = OpLoad %uint %50 Aligned 4 + OpStore %3 %51 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpBranch %53 + %53 = OpLabel + OpControlBarrier %uint_2 %uint_2 %uint_72 + %56 = OpAccessChain %_ptr_StorageBuffer_uint %21 %uint_0 %37 + %57 = OpLoad %uint %56 Aligned 4 + %58 = OpINotEqual %bool %57 %uint_0 + OpSelectionMerge %65 None + OpBranchConditional %58 %61 %65 + %61 = OpLabel + %62 = OpAtomicIAdd %uint %3 %uint_5 %uint_0 %uint_1 + %63 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %62 + OpStore %63 %37 Aligned 4 + OpBranch %65 + %65 = OpLabel + OpReturn + OpFunctionEnd + %81 = OpExtInst %void %66 PushConstantRegionOffset %uint_0 %uint_12 + %83 = OpExtInst %void %66 PushConstantRegionGroupOffset %uint_16 %uint_12 + %70 = OpExtInst %void %66 Kernel %26 %67 %uint_3 %uint_0 %68 + %72 = OpExtInst %void %66 ArgumentInfo %71 + %73 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_0 %uint_0 %uint_0 %72 + %75 = OpExtInst %void %66 ArgumentInfo %74 + %76 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_1 %uint_0 %uint_1 %75 + %78 = OpExtInst %void %66 ArgumentInfo %77 + %79 = OpExtInst %void %66 ArgumentStorageBuffer %70 %uint_2 %uint_0 %uint_2 %78 + %84 = OpExtInst %void %66 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %86 = OpExtInst %void %66 WorkgroupVariableSize %3 %uint_4 \ No newline at end of file diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-1.1.4.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-1.1.4.spvasm new file mode 100644 index 0000000000..00f5be77e4 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-1.1.4.spvasm @@ -0,0 +1,107 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 1, 1, 4 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 73 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %57 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %58 = OpString "histo_main_kernel" + %59 = OpString "__kernel" + %61 = OpString "sm_mappings" + %64 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_264 = OpConstant %uint 264 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %47 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %49 = OpINotEqual %bool %47 %uint_0 + OpSelectionMerge %56 None + OpBranchConditional %49 %52 %56 + %52 = OpLabel + %53 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %54 = OpAtomicIAdd %uint %53 %uint_5 %uint_0 %47 + OpBranch %56 + %56 = OpLabel + OpReturn + OpFunctionEnd + %68 = OpExtInst %void %57 PushConstantRegionOffset %uint_0 %uint_12 + %60 = OpExtInst %void %57 Kernel %26 %58 %uint_2 %uint_0 %59 + %62 = OpExtInst %void %57 ArgumentInfo %61 + %63 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_0 %uint_0 %uint_0 %62 + %65 = OpExtInst %void %57 ArgumentInfo %64 + %66 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_1 %uint_0 %uint_1 %65 + %69 = OpExtInst %void %57 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %71 = OpExtInst %void %57 WorkgroupVariableSize %5 %uint_8 diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-2.1.2.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-2.1.2.spvasm new file mode 100644 index 0000000000..f08b8e4495 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-2.1.2.spvasm @@ -0,0 +1,107 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 73 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %57 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %58 = OpString "histo_main_kernel" + %59 = OpString "__kernel" + %61 = OpString "sm_mappings" + %64 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_264 = OpConstant %uint 264 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %47 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %49 = OpINotEqual %bool %47 %uint_0 + OpSelectionMerge %56 None + OpBranchConditional %49 %52 %56 + %52 = OpLabel + %53 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %54 = OpAtomicIAdd %uint %53 %uint_5 %uint_0 %47 + OpBranch %56 + %56 = OpLabel + OpReturn + OpFunctionEnd + %68 = OpExtInst %void %57 PushConstantRegionOffset %uint_0 %uint_12 + %60 = OpExtInst %void %57 Kernel %26 %58 %uint_2 %uint_0 %59 + %62 = OpExtInst %void %57 ArgumentInfo %61 + %63 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_0 %uint_0 %uint_0 %62 + %65 = OpExtInst %void %57 ArgumentInfo %64 + %66 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_1 %uint_0 %uint_1 %65 + %69 = OpExtInst %void %57 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %71 = OpExtInst %void %57 WorkgroupVariableSize %5 %uint_8 diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-4.1.1.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-4.1.1.spvasm new file mode 100644 index 0000000000..4ba0820a40 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-4.1.1.spvasm @@ -0,0 +1,107 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 4, 1, 1 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 73 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %57 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %58 = OpString "histo_main_kernel" + %59 = OpString "__kernel" + %61 = OpString "sm_mappings" + %64 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_264 = OpConstant %uint 264 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %47 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %49 = OpINotEqual %bool %47 %uint_0 + OpSelectionMerge %56 None + OpBranchConditional %49 %52 %56 + %52 = OpLabel + %53 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %54 = OpAtomicIAdd %uint %53 %uint_5 %uint_0 %47 + OpBranch %56 + %56 = OpLabel + OpReturn + OpFunctionEnd + %68 = OpExtInst %void %57 PushConstantRegionOffset %uint_0 %uint_12 + %60 = OpExtInst %void %57 Kernel %26 %58 %uint_2 %uint_0 %59 + %62 = OpExtInst %void %57 ArgumentInfo %61 + %63 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_0 %uint_0 %uint_0 %62 + %65 = OpExtInst %void %57 ArgumentInfo %64 + %66 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_1 %uint_0 %uint_1 %65 + %69 = OpExtInst %void %57 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %71 = OpExtInst %void %57 WorkgroupVariableSize %5 %uint_8 diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-dv2wg.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-dv2wg.spvasm new file mode 100644 index 0000000000..f56dd6f668 --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-dv2wg.spvasm @@ -0,0 +1,106 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 72 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %57 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %58 = OpString "histo_main_kernel" + %59 = OpString "__kernel" + %61 = OpString "sm_mappings" + %64 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_264 = OpConstant %uint 264 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %47 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %49 = OpINotEqual %bool %47 %uint_0 + OpSelectionMerge %56 None + OpBranchConditional %49 %52 %56 + %52 = OpLabel + %53 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %54 = OpAtomicIAdd %uint %53 %uint_2 %uint_0 %47 + OpBranch %56 + %56 = OpLabel + OpReturn + OpFunctionEnd + %68 = OpExtInst %void %57 PushConstantRegionOffset %uint_0 %uint_12 + %60 = OpExtInst %void %57 Kernel %26 %58 %uint_2 %uint_0 %59 + %62 = OpExtInst %void %57 ArgumentInfo %61 + %63 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_0 %uint_0 %uint_0 %62 + %65 = OpExtInst %void %57 ArgumentInfo %64 + %66 = OpExtInst %void %57 ArgumentStorageBuffer %60 %uint_1 %uint_0 %uint_1 %65 + %69 = OpExtInst %void %57 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %71 = OpExtInst %void %57 WorkgroupVariableSize %5 %uint_8 diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-1.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-1.spvasm new file mode 100644 index 0000000000..4214872b9d --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-1.spvasm @@ -0,0 +1,108 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 74 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %58 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %59 = OpString "histo_main_kernel" + %60 = OpString "__kernel" + %62 = OpString "sm_mappings" + %65 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_72 = OpConstant %uint 72 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %uint_264 = OpConstant %uint 264 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_72 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %48 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %50 = OpINotEqual %bool %48 %uint_0 + OpSelectionMerge %57 None + OpBranchConditional %50 %53 %57 + %53 = OpLabel + %54 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %55 = OpAtomicIAdd %uint %54 %uint_5 %uint_0 %48 + OpBranch %57 + %57 = OpLabel + OpReturn + OpFunctionEnd + %69 = OpExtInst %void %58 PushConstantRegionOffset %uint_0 %uint_12 + %61 = OpExtInst %void %58 Kernel %26 %59 %uint_2 %uint_0 %60 + %63 = OpExtInst %void %58 ArgumentInfo %62 + %64 = OpExtInst %void %58 ArgumentStorageBuffer %61 %uint_0 %uint_0 %uint_0 %63 + %66 = OpExtInst %void %58 ArgumentInfo %65 + %67 = OpExtInst %void %58 ArgumentStorageBuffer %61 %uint_1 %uint_0 %uint_1 %66 + %70 = OpExtInst %void %58 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %72 = OpExtInst %void %58 WorkgroupVariableSize %5 %uint_8 diff --git a/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-2.spvasm b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-2.spvasm new file mode 100644 index 0000000000..5c407fc0aa --- /dev/null +++ b/dartagnan/src/test/resources/spirv/vulkan/ma/histogram-lc2gb-2.spvasm @@ -0,0 +1,108 @@ +; @Input: %22 = {{0, 1, 1, 1}} +; @Input: %23 = {{0, 0}} +; @Output: forall (%23[0][0] == 1 and %23[0][1] == 3) +; @Config: 2, 1, 2 +; SPIR-V +; Version: 1.6 +; Generator: Google Clspv; 0 +; Bound: 74 +; Schema: 0 + OpCapability Shader + OpCapability VulkanMemoryModel + OpExtension "SPV_KHR_vulkan_memory_model" + %58 = OpExtInstImport "NonSemantic.ClspvReflection.5" + OpMemoryModel Logical Vulkan + OpEntryPoint GLCompute %26 "histo_main_kernel" %5 %gl_GlobalInvocationID %gl_LocalInvocationID %18 %22 %23 %9 + OpSource OpenCL_C 200 + %59 = OpString "histo_main_kernel" + %60 = OpString "__kernel" + %62 = OpString "sm_mappings" + %65 = OpString "global_histo" + OpMemberDecorate %_struct_7 0 Offset 0 + OpDecorate %_struct_7 Block + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %_struct_20 0 Offset 0 + OpDecorate %_struct_20 Block + OpDecorate %22 DescriptorSet 0 + OpDecorate %22 Binding 0 + OpDecorate %23 DescriptorSet 0 + OpDecorate %23 Binding 1 + OpDecorate %13 SpecId 0 + OpDecorate %14 SpecId 1 + OpDecorate %15 SpecId 2 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_uint_uint_2 = OpTypeArray %uint %uint_2 +%_ptr_Workgroup__arr_uint_uint_2 = OpTypePointer Workgroup %_arr_uint_uint_2 + %v3uint = OpTypeVector %uint 3 + %_struct_7 = OpTypeStruct %v3uint +%_ptr_PushConstant__struct_7 = OpTypePointer PushConstant %_struct_7 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint + %13 = OpSpecConstant %uint 1 + %14 = OpSpecConstant %uint 1 + %15 = OpSpecConstant %uint 1 +%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %13 %14 %15 +%_ptr_Private_v3uint = OpTypePointer Private %v3uint +%_runtimearr_uint = OpTypeRuntimeArray %uint + %_struct_20 = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer__struct_20 = OpTypePointer StorageBuffer %_struct_20 + %void = OpTypeVoid + %25 = OpTypeFunction %void +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_PushConstant_uint = OpTypePointer PushConstant %uint +%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint + %uint_264 = OpConstant %uint 264 +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %uint_1 = OpConstant %uint 1 + %uint_72 = OpConstant %uint 72 + %bool = OpTypeBool + %uint_12 = OpConstant %uint 12 + %uint_8 = OpConstant %uint 8 + %5 = OpVariable %_ptr_Workgroup__arr_uint_uint_2 Workgroup + %9 = OpVariable %_ptr_PushConstant__struct_7 PushConstant +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input + %18 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize + %22 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %23 = OpVariable %_ptr_StorageBuffer__struct_20 StorageBuffer + %uint_5 = OpConstant %uint 5 + %26 = OpFunction %void None %25 + %27 = OpLabel + %30 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %31 = OpLoad %uint %30 Aligned 16 + %32 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %33 = OpLoad %uint %32 Aligned 16 + %35 = OpAccessChain %_ptr_PushConstant_uint %9 %uint_0 %uint_0 + %36 = OpLoad %uint %35 Aligned 16 + %38 = OpAccessChain %_ptr_Workgroup_uint %5 %31 + OpStore %38 %uint_0 Aligned|MakePointerAvailable|NonPrivatePointer 4 %uint_2 + OpControlBarrier %uint_2 %uint_2 %uint_264 + %40 = OpIAdd %uint %33 %36 + %42 = OpAccessChain %_ptr_StorageBuffer_uint %22 %uint_0 %40 + %43 = OpLoad %uint %42 Aligned 4 + %44 = OpAccessChain %_ptr_Workgroup_uint %5 %43 + %46 = OpAtomicIAdd %uint %44 %uint_2 %uint_0 %uint_1 + OpControlBarrier %uint_2 %uint_2 %uint_72 + %48 = OpLoad %uint %38 Aligned|MakePointerVisible|NonPrivatePointer 4 %uint_2 + %50 = OpINotEqual %bool %48 %uint_0 + OpSelectionMerge %57 None + OpBranchConditional %50 %53 %57 + %53 = OpLabel + %54 = OpAccessChain %_ptr_StorageBuffer_uint %23 %uint_0 %31 + %55 = OpAtomicIAdd %uint %54 %uint_5 %uint_0 %48 + OpBranch %57 + %57 = OpLabel + OpReturn + OpFunctionEnd + %69 = OpExtInst %void %58 PushConstantRegionOffset %uint_0 %uint_12 + %61 = OpExtInst %void %58 Kernel %26 %59 %uint_2 %uint_0 %60 + %63 = OpExtInst %void %58 ArgumentInfo %62 + %64 = OpExtInst %void %58 ArgumentStorageBuffer %61 %uint_0 %uint_0 %uint_0 %63 + %66 = OpExtInst %void %58 ArgumentInfo %65 + %67 = OpExtInst %void %58 ArgumentStorageBuffer %61 %uint_1 %uint_0 %uint_1 %66 + %70 = OpExtInst %void %58 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2 + %72 = OpExtInst %void %58 WorkgroupVariableSize %5 %uint_8 diff --git a/litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus b/litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus new file mode 100644 index 0000000000..520ee8348c --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/barrier-ordered.litmus @@ -0,0 +1,17 @@ +OPENCL barrier-ordered + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + atomic_store_explicit(x, 1, memory_order_relaxed); + B1: barrier(CLK_GLOBAL_MEM_FENCE); +} + +P1@wg 0, dev 0 (global int* x) { + B1: barrier(CLK_GLOBAL_MEM_FENCE); + int r2 = *x; +} + +forall (1:r2=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/barrier-unordered.litmus b/litmus/OPENCL/mixedAtomicity/barrier-unordered.litmus new file mode 100644 index 0000000000..6f9206aa28 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/barrier-unordered.litmus @@ -0,0 +1,17 @@ +OPENCL barrier-ordered + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + B1: barrier(CLK_GLOBAL_MEM_FENCE); + atomic_store_explicit(x, 1, memory_order_relaxed); +} + +P1@wg 0, dev 0 (global int* x) { + int r2 = *x; + B1: barrier(CLK_GLOBAL_MEM_FENCE); +} + +~exists (1:r2==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/co-ar.litmus b/litmus/OPENCL/mixedAtomicity/co-ar.litmus new file mode 100644 index 0000000000..6ece6328e2 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/co-ar.litmus @@ -0,0 +1,18 @@ +OPENCL barrier-ordered + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + atomic_store_explicit(x, 1, memory_order_relaxed); + B1: barrier(CLK_GLOBAL_MEM_FENCE); +} + +P1@wg 0, dev 0 (global int* x) { + *x = 2; + B1: barrier(CLK_GLOBAL_MEM_FENCE); + int r1 = atomic_load_explicit(x, memory_order_acquire, memory_scope_work_group); +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/co-nr.litmus b/litmus/OPENCL/mixedAtomicity/co-nr.litmus new file mode 100644 index 0000000000..99c4ab0075 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/co-nr.litmus @@ -0,0 +1,18 @@ +OPENCL barrier-ordered + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + atomic_store_explicit(x, 1, memory_order_relaxed); + B1: barrier(CLK_GLOBAL_MEM_FENCE); +} + +P1@wg 0, dev 0 (global int* x) { + *x = 2; + B1: barrier(CLK_GLOBAL_MEM_FENCE); + int r1 = *x; +} + +~exists (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/corr.litmus b/litmus/OPENCL/mixedAtomicity/corr.litmus new file mode 100644 index 0000000000..6f93915b5a --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/corr.litmus @@ -0,0 +1,16 @@ +OPENCL corr + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + atomic_store_explicit(x, 1, memory_order_release); +} + +P1@wg 0, dev 0 (global int* x) { + int r1 = atomic_load_explicit(x, memory_order_acquire, memory_scope_work_group); + int r2 = *x; +} + +~exists (1:r1==1 /\ 1:r2==0) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/fence-relacq.litmus b/litmus/OPENCL/mixedAtomicity/fence-relacq.litmus new file mode 100644 index 0000000000..8b8fd5c81d --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/fence-relacq.litmus @@ -0,0 +1,19 @@ +OPENCL MP + +{ + [x]=0; +} + +P0@wg 0, dev 0 (global int* x, global int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_device); + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_seq_cst,memory_scope_device); + atomic_store_explicit(y, 1, memory_order_relaxed, memory_scope_device); +} + +P1@wg 1, dev 0 (global int* x, global int* y) { + while (atomic_load_explicit(y, memory_order_relaxed, memory_scope_device) == 0) {} + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_seq_cst,memory_scope_device); + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/histogram.litmus b/litmus/OPENCL/mixedAtomicity/histogram.litmus new file mode 100644 index 0000000000..cff37e140e --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/histogram.litmus @@ -0,0 +1,41 @@ +OPENCL Histgram +// input: {0, 1, 1, 1} +{ +[local_histo_0] = 0; +[local_histo_1] = 0; +[local_histo_2] = 0; +[global_histo_0] = 0; +[global_histo_1] = 0; +} + +P0@wg 0, dev 0 (local int* local_histo_0, local int* local_histo_1, global int* global_histo_0, global int* global_histo_1) { + *local_histo_0 = 0; + *local_histo_1 = 0; + B1: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(local_histo_0, 1, memory_order_relaxed, memory_scope_work_group); + B2: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(global_histo_0, *local_histo_0, memory_order_relaxed, memory_scope_device); +} + +P1@wg 0, dev 0 (local int* local_histo_0, local int* local_histo_1, global int* global_histo_0, global int* global_histo_1) { + B1: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(local_histo_1, 1, memory_order_relaxed, memory_scope_work_group); + B2: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(global_histo_1, *local_histo_1, memory_order_relaxed, memory_scope_device); +} + +P2@wg 1, dev 0 (local int* local_histo_2, global int* global_histo_0, global int* global_histo_1) { + *local_histo_2 = 0; + B1: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(local_histo_2, 1, memory_order_relaxed, memory_scope_work_group); + B2: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(global_histo_1, *local_histo_2, memory_order_relaxed, memory_scope_device); +} + +P3@wg 1, dev 0 (local int* local_histo_2, global int* global_histo_0, global int* global_histo_1) { + B1: barrier(CLK_LOCAL_MEM_FENCE); + atomic_fetch_add_explicit(local_histo_2, 1, memory_order_relaxed, memory_scope_work_group); + B2: barrier(CLK_LOCAL_MEM_FENCE); +} + +forall (global_histo_0 = 1 /\ global_histo_1 = 3) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/lb-an.litmus b/litmus/OPENCL/mixedAtomicity/lb-an.litmus new file mode 100644 index 0000000000..7e126b7da9 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/lb-an.litmus @@ -0,0 +1,18 @@ +OPENCL LB + +{ +[x] = 0; +[y] = 0; +} + +P0@wg 0, dev 0 (global int* x, global int* y){ + if (*x == 1) + *y=1; +} + +P1@wg 0, dev 0 (global int* x, global int* y){ + if (*y == 1) + *x=1; +} + +~exists (x == 1 /\ y == 1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/lb-nn.litmus b/litmus/OPENCL/mixedAtomicity/lb-nn.litmus new file mode 100644 index 0000000000..02f3c856ac --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/lb-nn.litmus @@ -0,0 +1,18 @@ +OPENCL LB + +{ +[x] = 0; +[y] = 0; +} + +P0@wg 0, dev 0 (global int* x, global int* y){ + if (atomic_load_explicit(x, memory_order_relaxed, memory_scope_work_group) == 1) + *y=1; +} + +P1@wg 0, dev 0 (global int* x, global int* y){ + if (*y == 1) + *x=1; +} + +~exists (x == 1 /\ y == 1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/mp-an-relacq.litmus b/litmus/OPENCL/mixedAtomicity/mp-an-relacq.litmus new file mode 100644 index 0000000000..b9c7a04ebc --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/mp-an-relacq.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_release, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(atomic_load_explicit(y, memory_order_acquire, memory_scope_work_group) == 0) {} + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/mp-fence-rel-acq.litmus b/litmus/OPENCL/mixedAtomicity/mp-fence-rel-acq.litmus new file mode 100644 index 0000000000..875d6347cf --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/mp-fence-rel-acq.litmus @@ -0,0 +1,20 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_release,memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_relaxed, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(atomic_load_explicit(y, memory_order_relaxed, memory_scope_work_group) == 0) {} + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_acquire,memory_scope_work_group); + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/mp-na-relacq.litmus b/litmus/OPENCL/mixedAtomicity/mp-na-relacq.litmus new file mode 100644 index 0000000000..5df94b9c5f --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/mp-na-relacq.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + *x = 1; + atomic_store_explicit(y, 1, memory_order_release, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(atomic_load_explicit(y, memory_order_acquire, memory_scope_work_group) == 0) {} + int r1 = atomic_load_explicit(x, memory_order_relaxed, memory_scope_work_group); +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/partial-atomic.litmus b/litmus/OPENCL/mixedAtomicity/partial-atomic.litmus new file mode 100644 index 0000000000..97eab06359 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/partial-atomic.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; +} + +P0@wg 0, dev 0 (global int* x) { + *x = 1; + atomic_store_explicit(x, 2, memory_order_relaxed, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x) { + int r0 = *x; + int r1 = atomic_fetch_add_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + int r2 = *x; +} + +exists (1:r0=0 /\ 1:r1=0 /\ 1:r2=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/partial-atomic1.litmus b/litmus/OPENCL/mixedAtomicity/partial-atomic1.litmus new file mode 100644 index 0000000000..da2fc93919 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/partial-atomic1.litmus @@ -0,0 +1,15 @@ +OPENCL MP + +{ + [x]=0; +} + +P0@wg 0, dev 0 (global int* x) { + *x = 1; +} + +P1@wg 0, dev 0 (global int* x) { + int r0 = atomic_fetch_add_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); +} + +forall (1:r0=0 \/ 1:r0=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/partial-atomic2.litmus b/litmus/OPENCL/mixedAtomicity/partial-atomic2.litmus new file mode 100644 index 0000000000..c031101bf5 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/partial-atomic2.litmus @@ -0,0 +1,12 @@ +OPENCL MP + +{ + [x]=0; +} + +P0@wg 0, dev 0 (global int* x) { + *x = 1; + int r0 = atomic_load_explicit(x, memory_order_relaxed, memory_scope_work_group); +} + +forall (0:r0=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/relacq-fence-ordered.litmus b/litmus/OPENCL/mixedAtomicity/relacq-fence-ordered.litmus new file mode 100644 index 0000000000..89c2af1f8d --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/relacq-fence-ordered.litmus @@ -0,0 +1,20 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + *x = 1; + atomic_store_explicit(y, 1, memory_order_release, memory_scope_work_group); + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_release, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + int r0 = atomic_load_explicit(y, memory_order_acquire, memory_scope_work_group); + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE, memory_order_acquire, memory_scope_work_group); + int r1 = *x; +} + +exists (1:r0=1 /\ 1:r1=0) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/relrlx-mp-ordered.litmus b/litmus/OPENCL/mixedAtomicity/relrlx-mp-ordered.litmus new file mode 100644 index 0000000000..2d375ca94c --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/relrlx-mp-ordered.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_release, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global int* y) { + while(atomic_load_explicit(y, memory_order_relaxed, memory_scope_work_group) == 0) {} + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/rlxacq-mp-ordered.litmus b/litmus/OPENCL/mixedAtomicity/rlxacq-mp-ordered.litmus new file mode 100644 index 0000000000..a991186b51 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/rlxacq-mp-ordered.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_relaxed, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global int* y) { + while(atomic_load_explicit(y, memory_order_acquire, memory_scope_work_group) == 0) {} + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/rlxrlx-mp-ordered.litmus b/litmus/OPENCL/mixedAtomicity/rlxrlx-mp-ordered.litmus new file mode 100644 index 0000000000..91e0ef99d7 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/rlxrlx-mp-ordered.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_relaxed, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(atomic_load_explicit(y, memory_order_relaxed, memory_scope_work_group) == 0) {} + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/rmw-bar.litmus b/litmus/OPENCL/mixedAtomicity/rmw-bar.litmus new file mode 100644 index 0000000000..8bc91a4e1e --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/rmw-bar.litmus @@ -0,0 +1,22 @@ +OPENCL Histgram + +{ +[x] = 0; +[y] = 0; +} + +P0@wg 0, dev 0 (global int* x, global int* y) { + int r0 = atomic_fetch_add_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + B1: barrier(CLK_GLOBAL_MEM_FENCE); + int r1 = *x; + int r2 = atomic_fetch_add_explicit(y, r1, memory_order_relaxed, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global int* y) { + int r0 = atomic_fetch_add_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + B1: barrier(CLK_GLOBAL_MEM_FENCE); + int r1 = *x; + int r2 = atomic_fetch_add_explicit(y, r1, memory_order_relaxed, memory_scope_work_group); +} + +forall (y = 4) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/sc-fence-ordered.litmus b/litmus/OPENCL/mixedAtomicity/sc-fence-ordered.litmus new file mode 100644 index 0000000000..e976cf35dd --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/sc-fence-ordered.litmus @@ -0,0 +1,20 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + atomic_store_explicit(x, 1, memory_order_relaxed, memory_scope_work_group); + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_seq_cst,memory_scope_work_group); + atomic_store_explicit(y, 1, memory_order_release, memory_scope_work_group); +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(atomic_load_explicit(y, memory_order_relaxed, memory_scope_work_group) == 0) {} + atomic_work_item_fence(CLK_GLOBAL_MEM_FENCE,memory_order_seq_cst,memory_scope_work_group); + int r1 = *x; +} + +forall (1:r1=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/ttas-lock.litmus b/litmus/OPENCL/mixedAtomicity/ttas-lock.litmus new file mode 100644 index 0000000000..9b4a56499a --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/ttas-lock.litmus @@ -0,0 +1,19 @@ +OPENCL barrier-ordered + +{ +[x] = 0; +} + +P0@wg 0, dev 0 (global int* x) { + atomic_store_explicit(x, 1, memory_order_release); +} + +P1@wg 0, dev 0 (global int* x) { + int r1 = *x; + int r2 = 1; + if (r1 == 1) { + r2 = atomic_exchange_explicit(x, 0, memory_order_acq_rel); + } +} + +forall (1:r2=1) \ No newline at end of file diff --git a/litmus/OPENCL/mixedAtomicity/wkwk-mp-ordered.litmus b/litmus/OPENCL/mixedAtomicity/wkwk-mp-ordered.litmus new file mode 100644 index 0000000000..bcebd53e73 --- /dev/null +++ b/litmus/OPENCL/mixedAtomicity/wkwk-mp-ordered.litmus @@ -0,0 +1,18 @@ +OPENCL MP + +{ + [x]=0; + [y]=0; +} + +P0@wg 0, dev 0 (global int* x, global atomic_int* y) { + *x = 1; + *y = 1; +} + +P1@wg 0, dev 0 (global int* x, global atomic_int* y) { + while(*y == 0) {} + int r1 = *x; +} + +forall (1:r1==1) \ No newline at end of file diff --git a/litmus/OPENCL/overhauling/example7a.litmus b/litmus/OPENCL/overhauling/example7a.litmus index ad8a90c534..93fd889f87 100644 --- a/litmus/OPENCL/overhauling/example7a.litmus +++ b/litmus/OPENCL/overhauling/example7a.litmus @@ -15,7 +15,7 @@ P0@wg 0, dev 0 (global int* x, global int* y) { } } -P1@wg 0, dev 0 (global int* x, local atomic_int* y) { +P1@wg 0, dev 0 (global int* x, global int* y) { if (*y == 1) { *x = 1; } diff --git a/litmus/PTX/Manual/Mixed-Atomicity-barrier-ordered.litmus b/litmus/PTX/Manual/Mixed-Atomicity-barrier-ordered.litmus new file mode 100644 index 0000000000..fff0ed26f2 --- /dev/null +++ b/litmus/PTX/Manual/Mixed-Atomicity-barrier-ordered.litmus @@ -0,0 +1,12 @@ +PTX Mixed-Atomicity-barrier-ordered +{ +x=0; +P0:r1=0; +P1:r2=0; +} + P0@cta 0,gpu 0 | P1@cta 0,gpu 0 ; + st.weak x, 1 | ld.weak r1, x ; + bar.cta.sync 1 | bar.cta.sync 1 ; + | ld.relaxed.sys r2, x ; +forall +(P1:r2 == 1) \ No newline at end of file