Skip to content

Commit 445eebc

Browse files
badnikhilthewilsonator
badnikhil
authored and
thewilsonator
committed
DCompute: test comparison/equality hooks
1 parent e95c3b0 commit 445eebc

4 files changed

Lines changed: 183 additions & 0 deletions

File tree

‎tests/codegen/dcompute_comparison_hooks.d‎

Lines changed: 82 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,82 @@
1+
// Array `==` / `!=` / `<` / `>` / `<=` / `>=` in @compute device code lower to the
2+
// druntime hooks core.internal.array.equality.__equals and
3+
// core.internal.array.comparison.__cmp (plus helpers like isEqual/at). Those hooks
4+
// live in host-only druntime modules but are device-legal. Regression guard: they
5+
// used to be either rejected during dcompute semantic analysis or emitted as hollow
6+
// `declare`-only stubs. This test asserts the device IR contains real `define`d
7+
// bodies (never bare `declare`s) for every element type that actually instantiates
8+
// the templates.
9+
//
10+
// float / double / real / float[N] / struct-with-opEquals / struct-with-float-field
11+
// all bypass the int[] memcmp fast path, so the genuine __equals / __cmp templates
12+
// are instantiated and codegen'd for the device. (Plain POD structs and integral
13+
// element types take the inline-memcmp path instead -- see compilable/dcompute_comparison_hooks.d.)
14+
//
15+
// REQUIRES: target_NVPTX
16+
// RUN: %ldc -mdcompute-targets=cuda-700 -m64 -output-ll -output-o -c \
17+
// RUN: -mdcompute-file-prefix=comparison_hooks -Iinputs %s
18+
// RUN: FileCheck %s < comparison_hooks_cuda700_64.ll
19+
20+
@compute(CompileFor.deviceOnly) module dcompute_comparison_hooks;
21+
import ldc.dcompute;
22+
23+
struct SEq { int x; bool opEquals(ref const SEq o) const { return x == o.x; } }
24+
struct SFloat { float f; }
25+
struct SCmp { int x; int opCmp(ref const SCmp o) const { return x < o.x ? -1 : (x > o.x ? 1 : 0); } }
26+
27+
// ---- __equals (== / !=) ----------------------------------------------------
28+
29+
// Dynamic float[] equality -> __equals!(float,float).
30+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}__equalsTfTf
31+
@kernel void eq_float(float[] a, float[] b, bool* o) { *o = (a == b); }
32+
33+
// Dynamic double[] equality -> __equals!(double,double).
34+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}__equalsTdTd
35+
@kernel void eq_double(double[] a, double[] b, bool* o) { *o = (a == b); }
36+
37+
// Static float[4] equality -> __equals!(float,float, len, len).
38+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}__equalsTfTfVmi4
39+
@kernel void eq_static(float[4] a, float[4] b, bool* o) { *o = (a == b); }
40+
41+
// Struct with custom opEquals -> __equals!(SEq,SEq), element-wise via the user opEquals.
42+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}__equalsTS{{.*}}3SEq
43+
@kernel void eq_opequals(SEq[] a, SEq[] b, bool* o) { *o = (a == b); }
44+
45+
// Struct with a float field (forces element-wise compare, not memcmp) -> __equals!(SFloat,SFloat).
46+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}__equalsTS{{.*}}6SFloat
47+
@kernel void eq_floatfield(SFloat[] a, SFloat[] b, bool* o) { *o = (a == b); }
48+
49+
// `!=` lowers to the same __equals!(float,float) hook (negated). No NEW symbol needed
50+
// (shares eq_float's instantiation), but it must compile and stay defined.
51+
@kernel void ne_float(float[] a, float[] b, bool* o) { *o = (a != b); }
52+
53+
// The element comparator helper isEqual must also be defined (reached transitively
54+
// from __equals, lives in the same host-only equality module).
55+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array8equality{{.*}}isEqual
56+
57+
// ---- __cmp (< / > / <= / >=) -----------------------------------------------
58+
59+
// float[] ordering -> __cmp!float. All four relational operators lower to the
60+
// same __cmp!float hook, so one define covers <, >, <=, >=.
61+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array10comparison{{.*}}__cmpTf
62+
@kernel void cmp_lt(float[] a, float[] b, bool* o) { *o = (a < b); }
63+
@kernel void cmp_gt(float[] a, float[] b, bool* o) { *o = (a > b); }
64+
@kernel void cmp_le(float[] a, float[] b, bool* o) { *o = (a <= b); }
65+
@kernel void cmp_ge(float[] a, float[] b, bool* o) { *o = (a >= b); }
66+
67+
// double[] ordering -> __cmp!double.
68+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array10comparison{{.*}}__cmpTd
69+
@kernel void cmp_double(double[] a, double[] b, bool* o) { *o = (a < b); }
70+
71+
// int[] ordering -> __cmp!int (ordering goes through __cmp, NOT the memcmp fast path
72+
// that int[] `==` uses).
73+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array10comparison{{.*}}__cmpTi
74+
@kernel void cmp_int(int[] a, int[] b, bool* o) { *o = (a < b); }
75+
76+
// Struct with custom opCmp -> __cmp!(SCmp), element-wise via the user opCmp.
77+
// CHECK-DAG: define{{.*}}ptx_device{{.*}}@_D4core8internal5array10comparison{{.*}}__cmpTS{{.*}}4SCmp
78+
@kernel void cmp_opcmp(SCmp[] a, SCmp[] b, bool* o) { *o = (a < b); }
79+
80+
// None of the hooks may be left as hollow `declare`-only stubs.
81+
// CHECK-NOT: declare{{.*}}__equals
82+
// CHECK-NOT: declare{{.*}}__cmp

‎tests/compilable/dcompute_comparison_hooks.d‎

Lines changed: 53 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,53 @@
1+
// Companion to codegen/dcompute_comparison_hooks.d. That test proves the array
2+
// comparison/equality hooks emit defined device bodies; this one proves the
3+
// surrounding compile-ok matrix:
4+
// * the hooks compile in a deviceOnly @compute module, used directly in a
5+
// @kernel AND in a called non-kernel @compute helper,
6+
// * the hooks also compile in a hostAndDevice @compute module
7+
// (inputs/comparison_hooks_had.d),
8+
// * the Cat-2 "controls": integral element `==` takes the inline memcmp fast
9+
// path and never instantiates __equals at all.
10+
//
11+
// REQUIRES: target_NVPTX
12+
13+
// deviceOnly: hooks in a kernel and in a called @compute helper must compile.
14+
// RUN: %ldc -c -mdcompute-targets=cuda-700 %s
15+
16+
// hostAndDevice: the same hooks must compile when the module targets both.
17+
// RUN: %ldc -c -mdcompute-targets=cuda-700 -I%S %S/inputs/comparison_hooks_had.d
18+
19+
// Cat-2 control: int[]/int[4]/byte[] `==` lower to an inline memcmp; the __equals
20+
// template is NEVER instantiated. Assert the device IR has no __equals symbol but
21+
// does contain memcmp, documenting the bypass so future readers aren't surprised.
22+
// RUN: %ldc -mdcompute-targets=cuda-700 -m64 -output-ll -output-o -c \
23+
// RUN: -mdcompute-file-prefix=cmphooks_ctrl -d-version=Controls %s
24+
// RUN: FileCheck %s --check-prefix=CTRL < cmphooks_ctrl_cuda700_64.ll
25+
26+
@compute(CompileFor.deviceOnly) module dcompute_comparison_hooks;
27+
import ldc.dcompute;
28+
29+
struct SEq { int x; bool opEquals(ref const SEq o) const { return x == o.x; } }
30+
31+
version (Controls)
32+
{
33+
// CTRL-NOT: __equals
34+
// CTRL: memcmp
35+
@kernel void c_int (int[] a, int[] b, bool* o) { *o = (a == b); }
36+
@kernel void c_int4(int[4] a, int[4] b, bool* o) { *o = (a == b); }
37+
@kernel void c_byte(byte[] a, byte[] b, bool* o) { *o = (a == b); }
38+
}
39+
else
40+
{
41+
// Hook used in a non-kernel @compute helper that the kernel calls.
42+
bool eqHelper(float[] a, float[] b) { return a == b; }
43+
int cmpHelper(double[] a, double[] b) { return a < b ? -1 : 1; }
44+
45+
// Hook used directly in a kernel, across element types.
46+
@kernel void k_float (float[] a, float[] b, bool* o) { *o = (a == b); }
47+
@kernel void k_double(double[] a, double[] b, bool* o) { *o = (a < b); }
48+
@kernel void k_struct(SEq[] a, SEq[] b, bool* o) { *o = (a == b); }
49+
@kernel void k_helper(float[] a, float[] b, double[] c, double[] d, bool* o)
50+
{
51+
*o = eqHelper(a, b) && (cmpHelper(c, d) < 0);
52+
}
53+
}

‎tests/compilable/dcompute_ctfe_host_templates.d‎

Lines changed: 37 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,37 @@
1+
// No-regression guard (Cat 4): the comparison/equality-hook allowance must NOT
2+
// drag CTFE-only host templates into device legality checking.
3+
//
4+
// std.traits.fullyQualifiedName / moduleName are host-only templates whose
5+
// implementations use array concatenation (`~`) -- lowered to host runtime
6+
// helpers in core.internal.array.concatenation / .utils that do GC.malloc/memcpy
7+
// and are genuinely device-illegal. Used purely at compile time (here: to derive
8+
// CT lengths) they must not trip any device error, because the concat helpers run
9+
// only during CTFE and never reach device codegen. The hook allowance is scoped to
10+
// core.internal.array.comparison / .equality, so it must leave these alone.
11+
//
12+
// This is the exact case that regressed under an over-broad "any instantiated
13+
// template is device-legal" predicate; it must compile cleanly (exit 0).
14+
//
15+
// REQUIRES: target_NVPTX
16+
// RUN: %ldc -c -mdcompute-targets=cuda-700 %s
17+
18+
@compute(CompileFor.deviceOnly) module dcompute_ctfe_host_templates;
19+
import ldc.dcompute;
20+
import std.traits : fullyQualifiedName, moduleName;
21+
22+
struct SomeType { int x; }
23+
24+
@kernel void k(GlobalPointer!int o)
25+
{
26+
// CT-only consumption of the host template results -- never materializes the
27+
// string at runtime, so no host array machinery reaches device codegen.
28+
enum nFqnInt = fullyQualifiedName!int.length;
29+
enum nFqnType = fullyQualifiedName!SomeType.length;
30+
enum nModName = moduleName!SomeType.length;
31+
32+
int acc = 0;
33+
static foreach (i; 0 .. nFqnInt)
34+
acc += i;
35+
36+
*o = acc + cast(int)(nFqnType + nModName);
37+
}

‎tests/compilable/inputs/comparison_hooks_had.d‎

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
// hostAndDevice companion for compilable/dcompute_comparison_hooks.d.
2+
// The array comparison/equality hooks (__equals/__cmp) must also compile when the
3+
// @compute module targets both host and device, not just deviceOnly.
4+
@compute(CompileFor.hostAndDevice) module inputs.comparison_hooks_had;
5+
import ldc.dcompute;
6+
7+
struct SEq { int x; bool opEquals(ref const SEq o) const { return x == o.x; } }
8+
9+
@kernel void had_eq (float[] a, float[] b, bool* o) { *o = (a == b); }
10+
@kernel void had_cmp (float[] a, float[] b, bool* o) { *o = (a < b); }
11+
@kernel void had_struct(SEq[] a, SEq[] b, bool* o) { *o = (a == b); }

0 commit comments

Comments
 (0)