Support `constref` parameters passing. (#3249)

* Support `constref` parameters passing.

* Fix.

* Fix.

* Add test and diagnostic on mix use of __constref and no_diff.

* check for [constref] on differentiable member method.

---------

Co-authored-by: Yong He <yhe@nvidia.com>

5 files changed, 147 insertions, 4 deletions

diff --git a/tests/autodiff/constref-param.slang b/tests/autodiff/constref-param.slang
new file mode 100644
index 000000000..5bce8e304
--- /dev/null
+++ b/tests/autodiff/constref-param.slang
@@ -0,0 +1,41 @@
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-slang -compute -shaderobj -output-using-type
+//TEST(compute, vulkan):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-vk -compute -shaderobj -output-using-type
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-cuda -compute -shaderobj -output-using-type
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=CHECK):-cpu -compute -shaderobj -output-using-type
+
+//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
+RWStructuredBuffer<float> outputBuffer;
+
+struct NonDiff
+{
+    float a;
+}
+
+[Differentiable]
+float myFunc(__constref NonDiff fIn, float x, __constref no_diff float y)
+{
+    return x * fIn.a + y;
+}
+
+// test that the ordering of no_diff and __constref doesn't matter.
+[Differentiable]
+float myFunc2(__constref NonDiff fIn, float x, no_diff __constref float y)
+{
+    return x * fIn.a + y;
+}
+
+[numthreads(1, 1, 1)]
+void computeMain(uint3 dispatchThreadID: SV_DispatchThreadID)
+{
+    float a = 10.0;
+    NonDiff fIn = { a };
+    DifferentialPair<float> dpx = DifferentialPair<float>(4.f, 1.f);
+    float rs = __fwd_diff(myFunc)(fIn, dpx, 1.0).d;
+    float rs2 = __fwd_diff(myFunc2)(fIn, dpx, 1.0).d;
+
+    // CHECK: 10.0
+    // CHECK: 10.0
+
+    outputBuffer[0] = rs;
+    outputBuffer[1] = rs2;
+}
diff --git a/tests/diagnostics/const-ref-differentiable-param.slang b/tests/diagnostics/const-ref-differentiable-param.slang
new file mode 100644
index 000000000..c345826e7
--- /dev/null
+++ b/tests/diagnostics/const-ref-differentiable-param.slang
@@ -0,0 +1,38 @@
+//DIAGNOSTIC_TEST:SIMPLE(filecheck=CHECK):
+
+
+[Differentiable]
+float f(__constref float3 val)
+{
+    return val.x;
+}
+
+struct MyType : IDifferentiable
+{
+    // Error: cannot use constref on a differentiable member method of a differentiable type.
+    [Differentiable]
+    [constref] float compute(float x) { return 0; }
+
+    // OK
+    [Differentiable]
+    float compute1(float x) { return 0; }
+
+    // OK
+    [constref] float compute2(float x) { return 0;}
+}
+
+struct MyType2
+{
+    // OK.
+    [Differentiable]
+    [constref] float compute(float x) { return 0; }
+
+    // OK
+    [constref]
+    float compute1(float x) { return 0; }
+}
+
+// CHECK-DAG: {{.*}}(5): error 38034: cannot use '__constref' on a differentiable parameter.
+// CHECK-NOT {{.*}}error
+// CHECK-DAG: {{.*}}(14): error 38034: cannot use '[constref]' on a differentiable member method of a differentiable type.
+// CHECK-NOT {{.*}}error
diff --git a/tests/diagnostics/param-mutation.slang b/tests/diagnostics/param-mutation.slang
index 835e645c8..723c4c212 100644
--- a/tests/diagnostics/param-mutation.slang
+++ b/tests/diagnostics/param-mutation.slang
@@ -20,8 +20,6 @@ int doThing(MutatingStruct s, int v)
 
 // For non-copyable types (such as HitObject or NonCopyableStruct declared below), if passed as as `in`
 // should produce an error.
-// NOTE! This *doesn't* produce an error (or warning) because NonCopyable types are *implicitly* 
-// made *ref* when parsed as arguments. 
 
 [__NonCopyableType]
 struct NonCopyableStruct
@@ -32,7 +30,6 @@ struct NonCopyableStruct
 
 int doThing2(NonCopyableStruct s, int v)
 {
-    // Currently doesn't produce an error/warning because NonCopyableStruct is passed as *ref* implicitly.
     s.setValue(v + 1);
     return s.m_value;
 }
diff --git a/tests/diagnostics/param-mutation.slang.expected b/tests/diagnostics/param-mutation.slang.expected
index a1fb34fb7..314d6cbe3 100644
--- a/tests/diagnostics/param-mutation.slang.expected
+++ b/tests/diagnostics/param-mutation.slang.expected
@@ -1,8 +1,11 @@
-result code = 0
+result code = -1
 standard error = {
 tests/diagnostics/param-mutation.slang(17): warning 30068: mutating method 'setValue' called on `in` parameter 's'; changes will not be visible to caller. copy the parameter into a local variable if this behavior is intended
     s.setValue(v + 1);
               ^
+tests/diagnostics/param-mutation.slang(36): error 30067: mutating method 'setValue' called on `in` parameter 's'; changes will not be visible to caller. copy the parameter into a local variable if this behavior is intended
+    s.setValue(v + 1);
+              ^
 }
 standard output = {
 }
diff --git a/tests/language-feature/pointer/const-ref.slang b/tests/language-feature/pointer/const-ref.slang
new file mode 100644
index 000000000..f62fda697
--- /dev/null
+++ b/tests/language-feature/pointer/const-ref.slang
@@ -0,0 +1,64 @@
+// pointer-self-reference.slang
+
+//TEST:SIMPLE(filecheck=CHECK): -target cuda -entry computeMain -stage compute
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -slang -compute -output-using-type -shaderobj
+//TEST(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -vk -compute -output-using-type -shaderobj
+
+
+//TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer
+RWStructuredBuffer<int> outputBuffer;
+
+struct Thing
+{
+    int value;
+    int bigArray[128];
+
+    // Check that we are not inserting local variables that are copies of `this` parameter.
+    
+    // CHECK: __device__ int Thing_getSum{{.*}}(Thing{{.*}} * this{{.*}})
+    // CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}}
+    // CHECK: }
+    [constref]
+    int getSum()
+    {
+        int result = 0;
+        for (int i = 0; i < 128; i++)
+        {
+            result += bigArray[i];
+        }
+        return result;
+    }
+};
+
+// Check that we are not inserting local variables that are copies of `thing` parameter.
+
+// CHECK: __device__ int test{{.*}}(Thing{{.*}} * thing{{.*}})
+// CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}}
+// CHECK: }
+
+int test(__constref Thing thing)
+{
+    int sum0 = thing.getSum();
+    AllMemoryBarrier();
+    int sum1 = thing.getSum();
+    return sum0 + sum1;
+}
+
+int caller(Thing thing)
+{
+    return test(thing);
+}
+
+[numthreads(1, 1, 1)]
+void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
+{
+    int idx = dispatchThreadID.x;
+
+    Thing thing = {};
+
+    thing.bigArray[0] = 100;
+    thing.bigArray[99] = 1;
+
+    // BUFFER: 202
+    outputBuffer[idx] = caller(thing);
+}