Checkpoint commit of of GPU/CUDA/PTX/Sqlite3/TraceStore progress.

tpn · tpn · commit adcd28ca762e · 2017-05-13T19:44:47.000-04:00
diff --git a/Cu/Cu.cuh b/Cu/Cu.cuh
@@ -0,0 +1,81 @@
+/*++
+
+Copyright (c) 2017 Trent Nelson <trent@trent.me>
+
+Module Name:
+
+    Cu.cuh
+
+Abstract:
+
+    This module is the main header file for the Cu component.
+
+--*/
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+//
+// Define NT-style typedefs.
+//
+
+typedef char CHAR;
+typedef short SHORT;
+typedef long LONG;
+typedef wchar_t WCHAR;    // wc,   16-bit UNICODE character
+
+typedef WCHAR *PWCHAR, *LPWCH, *PWCH;
+
+typedef CHAR *PCHAR, *LPCH, *PCH;
+
+typedef float FLOAT;
+typedef double DOUBLE;
+typedef FLOAT *PFLOAT;
+typedef DOUBLE *PDOUBLE;
+
+typedef unsigned char UCHAR;
+typedef unsigned short USHORT;
+typedef unsigned long ULONG;
+
+typedef UCHAR *PUCHAR;
+typedef USHORT *PUSHORT;
+typedef ULONG *PULONG;
+
+typedef CHAR *PCHAR;
+typedef SHORT *PSHORT;
+typedef LONG *PLONG;
+
+typedef long long LONGLONG;
+typedef long long LONG64;
+typedef unsigned long long ULONGLONG;
+typedef unsigned long long ULONG64;
+
+typedef LONG64 *PLONG64;
+typedef ULONG64 *PULONG64;
+
+#define VOID    void
+
+//
+// Define CUDA macros and typedefs in NT style.
+//
+
+#define HOST __host__
+#define GLOBAL __global__
+#define DEVICE __device__
+#define GridDim gridDim
+#define BlockDim blockDim
+#define BlockIndex blockIdx
+#define ThreadIndex threadIdx
+
+
+#define FOR_EACH_1D(Index, Total)                           \
+    for (Index = BlockIndex.x * BlockDim.x + ThreadIndex.x; \
+         Index < Total;                                     \
+         Index += BlockDim.x * GridDim.x)
+
+#ifdef __cplusplus
+}
+#endif
+
+// vim:set ts=8 sw=4 sts=4 tw=80 expandtab syntax=cuda                         :
diff --git a/Cu/TraceStoreKernels.cu b/Cu/TraceStoreKernels.cu
@@ -16,14 +16,41 @@ Abstract:
 extern "C" {
 #endif
 
-__global__
-void saxpy(int n, float a, float *x, float *y)
+#include "Cu.cuh"
+#include <no_sal2.h>
+
+GLOBAL
+VOID
+SinglePrecisionAlphaXPlusY(
+    _In_ LONG Total,
+    _In_ FLOAT Alpha,
+    _In_ PFLOAT X,
+    _Out_ PFLOAT Y
+    )
 {
-    for (int i = blockIdx.x * blockDim.x + threadIdx.x;
-         i < n;
-         i += blockDim.x * gridDim.x) {
+    LONG Index;
 
-        y[i] = a * x[i] + y[i];
+    FOR_EACH_1D(Index, Total) {
+        Y[Index] = Alpha * X[Index] + Y[Index];
+    }
+}
+
+GLOBAL
+VOID
+DeltaTimestamp(
+    _In_ ULONG64 Total,
+    _In_ PULONG64 Timestamp,
+    _Out_ PULONG64 Delta
+    )
+{
+    ULONG64 Index;
+
+    if (ThreadIndex.x % 32 == 0) {
+        return;
+    }
+
+    FOR_EACH_1D(Index, Total) {
+        Delta[Index] = Timestamp[Index] - Timestamp[Index-1];
     }
 }
 
diff --git a/Rtl/Cu.h b/Rtl/Cu.h
@@ -483,7 +483,7 @@ typedef
 _Check_return_
 CU_RESULT
 (CU_DEVICE_GET)(
-    _Out_ PCU_DEVICE Device,
+    _Outptr_result_maybenull_ PCU_DEVICE Device,
     _In_opt_ LONG Ordinal
     );
 typedef CU_DEVICE_GET *PCU_DEVICE_GET;
@@ -624,7 +624,7 @@ CU_RESULT
     _In_z_ PCHAR Image,
     _In_ LONG NumberOfOptions,
     _In_reads_(NumberOfOptions) PCU_JIT_OPTION Options,
-    _Out_writes_(NumberOfOptions) PPCU_JIT_OPTION OptionValuesPointer
+    _Out_writes_(NumberOfOptions) PPVOID OptionValuesPointer
     );
 typedef CU_MODULE_LOAD_DATA_EX *PCU_MODULE_LOAD_DATA_EX;
 
diff --git a/Rtl/Rtl.c b/Rtl/Rtl.c
@@ -3001,6 +3001,48 @@ InitializeLargePages(PRTL Rtl)
     return TRUE;
 }
 
+PVOID
+TryMapViewOfFileNuma2(
+    PRTL Rtl,
+    HANDLE FileMappingHandle,
+    HANDLE ProcessHandle,
+    ULONG64 Offset,
+    PVOID BaseAddress,
+    SIZE_T ViewSize,
+    ULONG AllocationType,
+    ULONG PageProtection,
+    ULONG PreferredNode
+    )
+{
+    LARGE_INTEGER FileOffset;
+
+    if (!Rtl->MapViewOfFileNuma2) {
+        goto Fallback;
+    }
+
+    AllocationType = FilterLargePageFlags(Rtl, AllocationType);
+
+    return Rtl->MapViewOfFileNuma2(FileMappingHandle,
+                                   ProcessHandle,
+                                   Offset,
+                                   BaseAddress,
+                                   ViewSize,
+                                   AllocationType,
+                                   PageProtection,
+                                   PreferredNode);
+
+Fallback:
+
+    FileOffset.QuadPart = Offset;
+    return Rtl->MapViewOfFileExNuma(FileMappingHandle,
+                                    PageProtection,
+                                    FileOffset.HighPart,
+                                    FileOffset.LowPart,
+                                    ViewSize,
+                                    BaseAddress,
+                                    PreferredNode);
+}
+
 RTL_API PROBE_FOR_READ ProbeForRead;
 
 _Use_decl_annotations_
@@ -3143,6 +3185,8 @@ InitializeRtl(
         )
     );
 
+    Rtl->TryMapViewOfFileNuma2 = TryMapViewOfFileNuma2;
+
     Rtl->OutputDebugStringA = OutputDebugStringA;
     Rtl->OutputDebugStringW = OutputDebugStringW;
 
diff --git a/Rtl/Rtl.h b/Rtl/Rtl.h
@@ -233,6 +233,22 @@ extern "C" {
 #define RtlPointerToOffset(B,P)    ((ULONG_PTR)(((PCHAR)(P)) - ((PCHAR)(B))))
 #endif
 
+#ifndef FlagOn
+#define FlagOn(_F,_SF)        ((_F) & (_SF))
+#endif
+
+#ifndef BooleanFlagOn
+#define BooleanFlagOn(F,SF)   ((BOOLEAN)(((F) & (SF)) != 0))
+#endif
+
+#ifndef SetFlag
+#define SetFlag(_F,_SF)       ((_F) |= (_SF))
+#endif
+
+#ifndef ClearFlag
+#define ClearFlag(_F,_SF)     ((_F) &= ~(_SF))
+#endif
+
 //
 // Helper macros.
 //
@@ -2334,6 +2350,23 @@ PVOID
     );
 typedef MAP_VIEW_OF_FILE_NUMA2 *PMAP_VIEW_OF_FILE_NUMA2;
 
+typedef
+_Ret_maybenull_
+__out_data_source(FILE)
+PVOID
+(WINAPI TRY_MAP_VIEW_OF_FILE_NUMA2)(
+    _In_ struct _RTL *Rtl,
+    _In_ HANDLE FileMappingHandle,
+    _In_ HANDLE ProcessHandle,
+    _In_ ULONG64 Offset,
+    _In_opt_ PVOID BaseAddress,
+    _In_ SIZE_T ViewSize,
+    _In_ ULONG AllocationType,
+    _In_ ULONG PageProtection,
+    _In_ ULONG PreferredNode
+    );
+typedef TRY_MAP_VIEW_OF_FILE_NUMA2 *PTRY_MAP_VIEW_OF_FILE_NUMA2;
+
 typedef
 BOOL
 (WINAPI FLUSH_VIEW_OF_FILE)(
@@ -5846,6 +5879,10 @@ typedef struct _Struct_size_bytes_(SizeOfStruct) _RTL {
     PVIRTUAL_ALLOC TryLargePageVirtualAlloc;
     PVIRTUAL_ALLOC_EX TryLargePageVirtualAllocEx;
 
+    PMAP_VIEW_OF_FILE_EX_NUMA MapViewOfFileExNuma;
+    PMAP_VIEW_OF_FILE_NUMA2 MapViewOfFileNuma2;
+    PTRY_MAP_VIEW_OF_FILE_NUMA2 TryMapViewOfFileNuma2;
+
     PATEXIT atexit;
     PATEXITEX AtExitEx;
 
@@ -5877,9 +5914,6 @@ typedef struct _Struct_size_bytes_(SizeOfStruct) _RTL {
     PCRYPT_BINARY_TO_STRING_A CryptBinaryToStringA;
     PCRYPT_BINARY_TO_STRING_W CryptBinaryToStringW;
 
-    PMAP_VIEW_OF_FILE_EX_NUMA MapViewOfFileExNuma;
-    PMAP_VIEW_OF_FILE_NUMA2 MapViewOfFileNuma2;
-
     POUTPUT_DEBUG_STRING_A OutputDebugStringA;
     POUTPUT_DEBUG_STRING_W OutputDebugStringW;
 
@@ -5956,6 +5990,20 @@ typedef struct _Struct_size_bytes_(SizeOfStruct) _RTL {
 
 } RTL, *PRTL, **PPRTL;
 
+FORCEINLINE
+ULONG
+FilterLargePageFlags(
+    _In_ PRTL Rtl,
+    _In_ ULONG Flags
+    )
+{
+    if (!Rtl->Flags.IsLargePageEnabled) {
+        return Flags & ~(MEM_LARGE_PAGES | SEC_LARGE_PAGES);
+    } else {
+        return Flags;
+    }
+}
+
 FORCEINLINE
 ULONG
 TrailingZeros(
@@ -8482,6 +8530,56 @@ Return Value:
     );                                         \
 } while (0)
 
+/*++
+
+    VOID
+    READ_REG_DWORD(
+        _In_ HKEY Key,
+        _In_ LITERAL Name,
+        _In_ PDWORD DwordPointer
+        );
+
+Routine Description:
+
+    This is a helper macro for reading REG_DWORD values from the registry.
+
+Arguments:
+
+    Key - Supplies an HKEY handle that represents an open registry key with
+        appropriate read access.
+
+    Name - Name of the registry key to read.  This is converted into a literal
+        wide character string by the macro (e.g. MaxNoneRefCount will become
+        L"MaxNoneRefCount").
+
+    DwordPointer - Supplies a pointer to a DWORD that will receive the
+        registry key value.
+
+Return Value:
+
+    None.
+
+    N.B. If an error occurs, 0 will be written to DwordPointer.
+
+--*/
+
+#define READ_REG_DWORD(Key, Name, DwordPointer) do { \
+    ULONG DwordLength = sizeof(*DwordPointer);       \
+    Result = RegGetValueW(                           \
+        Key,                                         \
+        NULL,                                        \
+        L#Name,                                      \
+        RRF_RT_REG_DWORD,                            \
+        NULL,                                        \
+        (PVOID)DwordPointer,                         \
+        &DwordLength                                 \
+    );                                               \
+    if (Result != ERROR_SUCCESS) {                   \
+        *DwordPointer = 0;                           \
+    }                                                \
+} while (0)
+
+
 /*++
 
     VOID
diff --git a/TraceStore/TraceStore.vcxproj b/TraceStore/TraceStore.vcxproj
@@ -125,9 +125,6 @@
       <Project>{4868dda6-86d8-4fdf-ab71-8f9ea51a1142}</Project>
     </ProjectReference>
   </ItemGroup>
-  <ItemGroup>
-    <None Include="TraceStoreKernels.cu" />
-  </ItemGroup>
   <Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
   <ImportGroup Label="ExtensionTargets">
   </ImportGroup>
diff --git a/TraceStore/TraceStore.vcxproj.filters b/TraceStore/TraceStore.vcxproj.filters
@@ -144,11 +144,6 @@
       <Filter>Source Files</Filter>
     </ClCompile>
   </ItemGroup>
-  <ItemGroup>
-    <None Include="TraceStoreKernels.cu">
-      <Filter>Source Files</Filter>
-    </None>
-  </ItemGroup>
   <!--
   <ItemGroup>
     <MASM Include="..\Rtl\__chkstk.s" Condition="'$(Platform)' == 'x64'">
diff --git a/TraceStore/TraceStorePrivate.h b/TraceStore/TraceStorePrivate.h
diff --git a/TraceStore/TraceStoreSqlite3ExtInit.c b/TraceStore/TraceStoreSqlite3ExtInit.c
diff --git a/build-ptx.bat b/build-ptx.bat
