Fix various intptr_t issues by defining its width in `getIntTypeInfo` (#6786)

* Define a bit size for the intptr types

* Fix intptr_t sign

* Extend intptr test to check for previously broken operations

* Fix intptr vector test on CUDA

* Handle intptr size in getAnyValueSize

* Fix formatting

* Try with __ARM_ARCH_ISA_64

* On macs, int64_t != intptr_t

Yikes

* Move define to prelude header

* Also check apple in host-prelude

* Fix define location

9 files changed, 72 insertions, 6 deletions

diff --git a/include/slang.h b/include/slang.h
index e61c3cd7f..d65362762 100644
--- a/include/slang.h
+++ b/include/slang.h
@@ -422,7 +422,7 @@ convention for interface methods.
     #endif
 #elif defined(__arm__)
     #define SLANG_PROCESSOR_ARM 1
-#elif defined(_M_ARM64) || defined(__aarch64__)
+#elif defined(_M_ARM64) || defined(__aarch64__) || defined(__ARM_ARCH_ISA_A64)
     #define SLANG_PROCESSOR_ARM_64 1
 #elif defined(__EMSCRIPTEN__)
     #define SLANG_PROCESSOR_WASM 1
diff --git a/prelude/slang-cpp-host-prelude.h b/prelude/slang-cpp-host-prelude.h
index 8bc0f5cad..66f74052c 100644
--- a/prelude/slang-cpp-host-prelude.h
+++ b/prelude/slang-cpp-host-prelude.h
@@ -28,6 +28,13 @@
 #include <string.h>
 #endif // SLANG_LLVM
 
+// Is intptr_t not equal to equal-width sized integer type?
+#if defined(__APPLE__)
+#define SLANG_INTPTR_TYPE_IS_DISTINCT 1
+#else
+#define SLANG_INTPTR_TYPE_IS_DISTINCT 0
+#endif
+
 #if defined(_MSC_VER)
 #define SLANG_PRELUDE_SHARED_LIB_EXPORT __declspec(dllexport)
 #else
diff --git a/prelude/slang-cpp-prelude.h b/prelude/slang-cpp-prelude.h
index 4dacac9c5..6530654c1 100644
--- a/prelude/slang-cpp-prelude.h
+++ b/prelude/slang-cpp-prelude.h
@@ -22,6 +22,13 @@
 #include <string.h>
 #endif // SLANG_LLVM
 
+// Is intptr_t not equal to equal-width sized integer type?
+#if defined(__APPLE__)
+#define SLANG_INTPTR_TYPE_IS_DISTINCT 1
+#else
+#define SLANG_INTPTR_TYPE_IS_DISTINCT 0
+#endif
+
 #if defined(_MSC_VER)
 #define SLANG_PRELUDE_SHARED_LIB_EXPORT __declspec(dllexport)
 #else
diff --git a/prelude/slang-cpp-types-core.h b/prelude/slang-cpp-types-core.h
index 82674fefb..5f2024ac7 100644
--- a/prelude/slang-cpp-types-core.h
+++ b/prelude/slang-cpp-types-core.h
@@ -346,6 +346,10 @@ SLANG_INT_VECTOR_OPS(uint)
 SLANG_INT_VECTOR_OPS(uint8_t)
 SLANG_INT_VECTOR_OPS(uint16_t)
 SLANG_INT_VECTOR_OPS(uint64_t)
+#if SLANG_INTPTR_TYPE_IS_DISTINCT
+SLANG_INT_VECTOR_OPS(intptr_t)
+SLANG_INT_VECTOR_OPS(uintptr_t)
+#endif
 
 SLANG_FLOAT_VECTOR_OPS(float)
 SLANG_FLOAT_VECTOR_OPS(double)
@@ -367,6 +371,10 @@ SLANG_VECTOR_INT_NEG_OP(uint)
 SLANG_VECTOR_INT_NEG_OP(uint8_t)
 SLANG_VECTOR_INT_NEG_OP(uint16_t)
 SLANG_VECTOR_INT_NEG_OP(uint64_t)
+#if SLANG_INTPTR_TYPE_IS_DISTINCT
+SLANG_VECTOR_INT_NEG_OP(intptr_t)
+SLANG_VECTOR_INT_NEG_OP(uintptr_t)
+#endif
 
 #define SLANG_FLOAT_VECTOR_MOD(T)                                               \
     template<int N>                                                             \
@@ -643,6 +651,10 @@ SLANG_INT_MATRIX_OPS(uint)
 SLANG_INT_MATRIX_OPS(uint8_t)
 SLANG_INT_MATRIX_OPS(uint16_t)
 SLANG_INT_MATRIX_OPS(uint64_t)
+#if SLANG_INTPTR_TYPE_IS_DISTINCT
+SLANG_INT_MATRIX_OPS(intptr_t)
+SLANG_INT_MATRIX_OPS(uintptr_t)
+#endif
 
 SLANG_FLOAT_MATRIX_OPS(float)
 SLANG_FLOAT_MATRIX_OPS(double)
@@ -665,6 +677,10 @@ SLANG_MATRIX_INT_NEG_OP(uint)
 SLANG_MATRIX_INT_NEG_OP(uint8_t)
 SLANG_MATRIX_INT_NEG_OP(uint16_t)
 SLANG_MATRIX_INT_NEG_OP(uint64_t)
+#if SLANG_INTPTR_TYPE_IS_DISTINCT
+SLANG_MATRIX_INT_NEG_OP(intptr_t)
+SLANG_MATRIX_INT_NEG_OP(uintptr_t)
+#endif
 
 #define SLANG_FLOAT_MATRIX_MOD(T)                                                             \
     template<int R, int C>                                                                    \
diff --git a/source/slang/slang-emit-cuda.cpp b/source/slang/slang-emit-cuda.cpp
index 8657b3707..bb6941907 100644
--- a/source/slang/slang-emit-cuda.cpp
+++ b/source/slang/slang-emit-cuda.cpp
@@ -123,6 +123,18 @@ UnownedStringSlice CUDASourceEmitter::getVectorPrefix(IROp op)
     case kIROp_UInt64Type:
         return UnownedStringSlice("ulonglong");
 
+#if SLANG_PTR_IS_64
+    case kIROp_IntPtrType:
+        return UnownedStringSlice("longlong");
+    case kIROp_UIntPtrType:
+        return UnownedStringSlice("ulonglong");
+#else
+    case kIROp_IntPtrType:
+        return UnownedStringSlice("int");
+    case kIROp_UIntPtrType:
+        return UnownedStringSlice("uint");
+#endif
+
     case kIROp_HalfType:
         return UnownedStringSlice("__half");
 
diff --git a/source/slang/slang-ir-any-value-marshalling.cpp b/source/slang/slang-ir-any-value-marshalling.cpp
index b3bcf3316..9fb414db1 100644
--- a/source/slang/slang-ir-any-value-marshalling.cpp
+++ b/source/slang/slang-ir-any-value-marshalling.cpp
@@ -862,11 +862,19 @@ SlangInt _getAnyValueSizeRaw(IRType* type, SlangInt offset)
     case kIROp_FloatType:
     case kIROp_UIntType:
     case kIROp_BoolType:
+#if SLANG_PTR_IS_32
+    case kIROp_IntPtrType:
+    case kIROp_UIntPtrType:
+#endif
         return alignUp(offset, 4) + 4;
     case kIROp_UInt64Type:
     case kIROp_Int64Type:
     case kIROp_DoubleType:
     case kIROp_PtrType:
+#if SLANG_PTR_IS_64
+    case kIROp_IntPtrType:
+    case kIROp_UIntPtrType:
+#endif
         return alignUp(offset, 8) + 8;
     case kIROp_Int16Type:
     case kIROp_UInt16Type:
diff --git a/source/slang/slang-ir.cpp b/source/slang/slang-ir.cpp
index c1cec36fe..a99eddebb 100644
--- a/source/slang/slang-ir.cpp
+++ b/source/slang/slang-ir.cpp
@@ -7769,6 +7769,12 @@ IntInfo getIntTypeInfo(const IRType* intType)
         return {32, false};
     case kIROp_UInt64Type:
         return {64, false};
+    case kIROp_UIntPtrType:
+#if SLANG_PTR_IS_32
+        return {32, false};
+#else
+        return {64, false};
+#endif
     case kIROp_Int8Type:
         return {8, true};
     case kIROp_Int16Type:
@@ -7777,9 +7783,12 @@ IntInfo getIntTypeInfo(const IRType* intType)
         return {32, true};
     case kIROp_Int64Type:
         return {64, true};
-
-    case kIROp_IntPtrType:  // target platform dependent
-    case kIROp_UIntPtrType: // target platform dependent
+    case kIROp_IntPtrType:
+#if SLANG_PTR_IS_32
+        return {32, true};
+#else
+        return {64, true};
+#endif
     default:
         SLANG_UNEXPECTED("Unhandled type passed to getIntTypeInfo");
     }
diff --git a/tests/language-feature/types/intptr.slang b/tests/language-feature/types/intptr.slang
index 0be67cdb7..491fc6c38 100644
--- a/tests/language-feature/types/intptr.slang
+++ b/tests/language-feature/types/intptr.slang
@@ -29,11 +29,16 @@ void computeMain(int3 dispatchThreadID: SV_DispatchThreadID)
     {
         testResult = 0;
     }
-    uintptr_t b = 0uz;
+    uintptr_t b = 50uz;
     actualSize = getSizeOf(b);
     if (expectedSize != actualSize)
     {
         testResult = 0;
     }
     outputBuffer[dispatchThreadID.x] = testResult;
+    outputBuffer[dispatchThreadID.x+1] = int(b >> 1);
+
+    vector<uintptr_t, 2> c = vector<uintptr_t, 2>(2uz);
+    vector<uintptr_t, 2> d = vector<uintptr_t, 2>(3uz);
+    outputBuffer[dispatchThreadID.x+2] = int((c+d).y);
 }
diff --git a/tests/language-feature/types/intptr.slang.expected.txt b/tests/language-feature/types/intptr.slang.expected.txt
index 56a6051ca..c78bed2b0 100644
--- a/tests/language-feature/types/intptr.slang.expected.txt
+++ b/tests/language-feature/types/intptr.slang.expected.txt
@@ -1 +1,3 @@
-1
\ No newline at end of file
+1
+19
+5