diff options
| author | ArielG-NV <159081215+ArielG-NV@users.noreply.github.com> | 2025-08-07 00:22:22 -0700 |
|---|---|---|
| committer | GitHub <noreply@github.com> | 2025-08-07 07:22:22 +0000 |
| commit | 063cbeaaea2fb00a10c6058ea4a9632092772ea5 (patch) | |
| tree | b4412347d6c264c3b1a84ec971921a5e2fe76134 /tests/language-feature | |
| parent | 9e2685853033f4286feaf22d04a755a7395d95ce (diff) | |
Initial copy elision pass (#8042)
Fixes #7574
Changes:
* Add an initial (fairly simple) optimization pass which is able to
eliminate redundant copies.
* Our current existing optimizer passes remove redundant load/store very
robustly, this pass will focus on other cases of copy elimination
* Primary approach is to make all functions which are `in T` and `T` is
trivial to copy into a `__constref T`. We then (depending on scenario)
manually insert a variable+load if a pass-by-reference is not possible;
otherwise we pass by `constref`.
* Added optimizations to eliminate redundant code which causes
`constref` to fail to compile
---------
Co-authored-by: Harsh Aggarwal <haaggarwal@nvidia.com>
Co-authored-by: Claude <noreply@anthropic.com>
Co-authored-by: slangbot <ellieh+slangbot@nvidia.com>
Co-authored-by: slangbot <186143334+slangbot@users.noreply.github.com>
Diffstat (limited to 'tests/language-feature')
| -rw-r--r-- | tests/language-feature/pointer/const-ref.slang | 8 |
1 files changed, 4 insertions, 4 deletions
diff --git a/tests/language-feature/pointer/const-ref.slang b/tests/language-feature/pointer/const-ref.slang index f62fda697..06bb9dc07 100644 --- a/tests/language-feature/pointer/const-ref.slang +++ b/tests/language-feature/pointer/const-ref.slang @@ -3,7 +3,7 @@ //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(compute):COMPARE_COMPUTE_EX(filecheck-buffer=BUFFER): -cuda -compute -output-using-type -shaderobj //TEST_INPUT:ubuffer(data=[0 0 0 0], stride=4):out,name=outputBuffer RWStructuredBuffer<int> outputBuffer; @@ -14,8 +14,8 @@ struct Thing 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: __device__ int Thing_getSum{{.*}}Thing{{.*}}*{{.*}}this{{.*}}) // CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}} // CHECK: } [constref] @@ -32,7 +32,7 @@ struct Thing // Check that we are not inserting local variables that are copies of `thing` parameter. -// CHECK: __device__ int test{{.*}}(Thing{{.*}} * thing{{.*}}) +// CHECK: __device__ int test{{.*}}Thing{{.*}}*{{.*}}thing{{.*}}) // CHECK-NOT: Thing{{[a-zA-Z0-9_]*}} {{[a-zA-Z0-9_]+}} // CHECK: } |