add const cache linesize detect

sjfeng1999 · sjfeng1999 · commit 00f391617e38 · 2022-03-28T23:23:55.000+08:00
diff --git a/README.md b/README.md
@@ -12,39 +12,49 @@
 2. `cmake .. && make`
 3. `python ../compile_sass.py -arch=<70|75|80>`
 
-## Benchmark
+## Microbenchmark
 
 ### 1. Latency
 
-|Device                      |           |  RTX-2070 |
-|:--------------------------:|:---------:|:---------:|
-|Global Latency              |cycle      | TBD       |
-|L2 Latency                  |cycle      | 236       |
-|L1 Latency                  |cycle      | 32        |  
-|Shared Latency              |cycle      | 23        |  
-|Constant Latency            |cycle      | 448       |
-|Constant L2 Latency         |cycle      | 62        |
-|Constant L1 Latency         |cycle      | 4         |  
+|Device                      |           |  Turing RTX-2070 |
+|:--------------------------:|:---------:|:----------------:|
+|Global Latency              |cycle      | TBD              |
+|L2 Latency                  |cycle      | 236              |
+|L1 Latency                  |cycle      | 32               |  
+|Shared Latency              |cycle      | 23               |  
+|Constant Latency            |cycle      | 448              |
+|Constant L2 Latency         |cycle      | 62               |
+|Constant L1 Latency         |cycle      | 4                |  
 
 - const L1-cache is as fast as register.
 
 ### 2. Cache Linesize
 
 
-|Device                      |           |  RTX-2070 |
-|:--------------------------:|:---------:|:---------:|
-|L2 Linesise                 |bytes      | TBD       |
-|L1 Linesize                 |bytes      | 32        |
-|Constant L2 Linesise        |bytes      | TBD       |
-|Constant L1 Linesize        |bytes      | TBD       |
+|Device                      |           | Turing RTX-2070  |
+|:--------------------------:|:---------:|:----------------:|
+|L2 Linesise                 |bytes      | 64               |
+|L1 Linesize                 |bytes      | 32               |
+|Constant L2 Linesise        |bytes      | 256              |
+|Constant L1 Linesize        |bytes      | 32               |
 
 
 
 ### 3. Reg Bankconflict
 
-| Instruction | conflict | without conflict |
-|:-----------:|:--------:|:----------------:|
-|FFMA         | 1.484    | 1.758            |
+| Instruction |         | conflict | without conflict |
+|:-----------:|:-------:|:--------:|:----------------:|
+|FFMA         |  CPI    | 1.758    | 1.484            |
+
+
+### 4. Shared Bankconflict
+
+| Memory Load            |           | Turing RTX-2070  |
+|:----------------------:|:---------:|:----------------:|
+| Single                 | cycle     |  23              |
+| Vector2 X 2            | cycle     |  27              |
+| Conflict Strided       | cycle     |  41              |
+| Conlict-Free Strided   | cycle     |  32              |
 
 
 
diff --git a/memory/cache_linesize.cu b/memory/cache_linesize.cu
@@ -7,16 +7,39 @@
 #include "utils.cuh"
 
 
-__global__ void linesizeDetectKernel(float* input, float* output, uint* clock){
+__constant__ float cinput[1024];
 
-    uint c[48];
+__global__ void linesizeDetectKernel(float* input, float* output, uint* clock, float* cinput){
+
+    uint c[256];
     float val = 0;
-    float acc = 0;
 
+    float acc = 0;
     c[0] = getClock();
+    #pragma unroll
+    for (int i = 0; i < 256; ++i){
+        asm volatile(
+            "ld.global.cg.b32    %0,    [%1];  \n\t"
+            :"=f"(val):"l"(input):"memory"
+        );
+        c[i+1] = getClock();
+        acc += val;
+        input += 2;
+    }
+    #pragma unroll
+    for (int i = 0; i < 256; ++i){
+        clock[i] = c[i+1] - c[i];
+    }
+    output[0] = acc;
+
+    /////////////////////////////////////////////////////////////////////////
 
+    input += 1024;
+    clock += 512;
+    acc = 0;
+    c[0] = getClock();
     #pragma unroll
-    for (int i = 0; i < 32; ++i){
+    for (int i = 0; i < 256; ++i){
         asm volatile(
             "ld.global.ca.f32    %0,    [%1];  \n\t"
             :"=f"(val):"l"(input):"memory"
@@ -25,27 +48,34 @@ __global__ void linesizeDetectKernel(float* input, float* output, uint* clock){
         acc += val;
         input++;
     }
-
     #pragma unroll
-    for (int i = 0; i < 32; ++i){
+    for (int i = 0; i < 256; ++i){
         clock[i] = c[i+1] - c[i];
     }
-    output[0] = acc;
+    output[1] = acc;
 }
 
 
-int detectCacheLinesize(uint* clock, int size){
-    int l1_linesize = 0;
+int detectCacheLinesize(uint* clock, int size, uint gap){
+    int linesize = 0;
     uint last_cycle = clock[0];
+
+    int first = 0;
+    int second = 0;
+
+    // formatArray(clock, 256, 16);
     for (int i = 1; i < size; ++i){
-        // printf("clock %d   latency %u\n", i, clock[i]);
-        if (clock[i] > last_cycle and clock[i] - last_cycle > 10) {
-            l1_linesize = i * 4;
-            break;
+        if (clock[i] > last_cycle and clock[i] - last_cycle > gap) {
+            if (first == 0){
+                first = i;
+            } else {
+                second = i;
+                break;
+            }
         } 
         last_cycle = clock[i];
     }
-    return l1_linesize;
+    return (second - first) * 4;
 }
 
 
@@ -72,25 +102,27 @@ int main(){
     dim3 gDim(1, 1, 1);
     dim3 bDim(1, 1, 1);
 
-    void* kernel_args[] = {&input_d, &output_d, &clock_d};
+    void* kernel_args[] = {&input_d, &output_d, &clock_d, &cinput};
     const char* cubin_name = "../sass_cubin/cache_linesize.cubin";
     const char* kernel_name = "cacheLinesize";
 
     launchSassKernel(cubin_name, kernel_name, gDim, bDim, 0, kernel_args);
     cudaMemcpy(clock_h, clock_d, sizeof(float) * size, cudaMemcpyDeviceToHost);
     cudaDeviceSynchronize();
     printf(">>> SASS Level Cache Linesize Result\n");
-    // printf("        L2 Linesize \t= %3d B\n", clock_h[1]);
-    printf("        L1 LineSize \t= %3u B\n", detectCacheLinesize(clock_h, size));
+    printf("        Global   L2 LineSize \t= %3u B\n", detectCacheLinesize(clock_h, 512, 40));
+    printf("        Global   L1 LineSize \t= %3u B\n", detectCacheLinesize(clock_h + 512,  512, 10));
+    printf("        Constant L2 LineSize \t= %3u B\n", detectCacheLinesize(clock_h + 1024, 512, 100));
+    printf("        Constant L1 LineSize \t= %3u B\n", detectCacheLinesize(clock_h + 1536, 512, 10));
 
 
 
-    linesizeDetectKernel<<<gDim, bDim>>>(input_d, output_d, clock_d);
+    linesizeDetectKernel<<<gDim, bDim>>>(input_d, output_d, clock_d, cinput);
     cudaMemcpy(clock_h, clock_d, sizeof(float) * size, cudaMemcpyDeviceToHost);
     cudaDeviceSynchronize();
     printf("\n");
     printf(">>> CUDA-C Level Cache Linesize Result\n");
-    // printf("        L2 Linesize \t= %3d B\n", clock_h[1]);
-    printf("        L1 LineSize \t= %3u B\n", detectCacheLinesize(clock_h, size));
+    printf("        Global   L2 LineSize \t= %3u B\n", detectCacheLinesize(clock_h, 512, 40));
+    printf("        Global   L1 LineSize \t= %3u B\n", detectCacheLinesize(clock_h + 512, 512, 10));
     return 0;
 }
diff --git a/sass_cubin/cache_linesize.sass b/sass_cubin/cache_linesize.sass
@@ -11,40 +11,118 @@ clock, 8
 3: output_hi
 4: clock_lo
 5: clock_hi
-6-48 ~ e<0-35>
-49: vA
-50-96 ~ c<0-35>
+6: vA
+7: e
+8-240 ~ c<0-200>
 </regs>
 
+
+<consts>
+const_a, 1024
+</consts>
+
+
+
 --:-:-:-:2      MOV input_lo, input[0];
 --:-:-:-:2      MOV input_hi, input[1];
 --:-:-:-:2      MOV output_lo, output[0];
 --:-:-:-:2      MOV output_hi, output[1];
 --:-:-:-:2      MOV clock_lo, clock[0];
 --:-:-:-:4      MOV clock_hi, clock[1];
 
---:-:-:-:2      CS2R c0, SR_CLOCKLO;
 
+--:-:-:-:2      CS2R c0, SR_CLOCKLO;
 <CODE>
 SASS_CODE = []
+loop_size = 200
+
+LDG  = "--:-:0:-:2      LDG.E.STRONG.GPU vA, [input_lo+{:}];"
+CS2R = "01:-:-:-:4      CS2R c{:}, SR_CLOCKLO;"
+IADD = "--:-:-:-:5      IADD3 e, c{:}, -c{:}, RZ;"
+STG  = "--:-:-:-:4      STG.E.SYS [clock_lo+{:}], e;"
+
+for i in range(loop_size):
+    SASS_CODE += [LDG.format(hex(i * 4))]
+    SASS_CODE += [CS2R.format(i+1)]
+
+for i in range(loop_size):
+    SASS_CODE += [IADD.format(i+1, i)]
+    SASS_CODE += [STG.format(hex(i*4), i)]
+
+out_ = "\n" + "\n".join(SASS_CODE) + "\n"
+</CODE>
+
 
-loop_size = 35
+
+
+
+
+--:-:-:-:2      IADD3 clock_lo, clock_lo, 0x800, RZ;
+--:-:-:-:2      CS2R c0, SR_CLOCKLO;
+
+<CODE>
+SASS_CODE = []
+loop_size = 200
 
 LDG  = "--:-:0:-:2      LDG.E.STRONG.CTA vA, [input_lo+{:}];"
 CS2R = "01:-:-:-:4      CS2R c{:}, SR_CLOCKLO;"
-IADD = "--:-:-:-:5      IADD3 e{:}, c{:}, -c{:}, RZ;"
-STG  = "--:-:-:-:4      STG.E.SYS [clock_lo+{:}], e{:};"
-BAR  = "--:-:-:-:2      BAR.SYNC 0x0;"
+IADD = "--:-:-:-:5      IADD3 e, c{:}, -c{:}, RZ;"
+STG  = "--:-:-:-:4      STG.E.SYS [clock_lo+{:}], e;"
 
 for i in range(loop_size):
     SASS_CODE += [LDG.format(hex(i * 4))]
     SASS_CODE += [CS2R.format(i+1)]
 
-SASS_CODE += [BAR]
 for i in range(loop_size):
-    SASS_CODE += [IADD.format(i, i+1, i)]
+    SASS_CODE += [IADD.format(i+1, i)]
+    SASS_CODE += [STG.format(hex(i*4), i)]
+
+out_ = "\n" + "\n".join(SASS_CODE) + "\n"
+</CODE>
+
+
+
+
+--:-:-:-:2      IADD3 clock_lo, clock_lo, 0x800, RZ;
+--:-:-:-:2      CS2R c0, SR_CLOCKLO;
+<CODE>
+SASS_CODE = []
+loop_size = 200
+
+LDC  = "--:-:0:-:2      LDC.E vA, const_a[{:}];"
+CS2R = "01:-:-:-:4      CS2R c{:}, SR_CLOCKLO;"
+IADD = "--:-:-:-:5      IADD3 e, c{:}, -c{:}, RZ;"
+STG  = "--:-:-:-:4      STG.E.SYS [clock_lo+{:}], e;"
+
+for i in range(loop_size):
+    SASS_CODE += [LDC.format(i)]
+    SASS_CODE += [CS2R.format(i+1)]
+
+for i in range(loop_size):
+    SASS_CODE += [IADD.format(i+1, i)]
+    SASS_CODE += [STG.format(hex(i*4), i)]
+
+out_ = "\n" + "\n".join(SASS_CODE) + "\n"
+</CODE>
+
+
+--:-:-:-:2      IADD3 clock_lo, clock_lo, 0x800, RZ;
+--:-:-:-:2      CS2R c0, SR_CLOCKLO;
+<CODE>
+SASS_CODE = []
+loop_size = 200
+
+LDC  = "--:-:0:-:2      MOV vA, const_a[{:}];"
+CS2R = "01:-:-:-:4      CS2R c{:}, SR_CLOCKLO;"
+IADD = "--:-:-:-:5      IADD3 e, c{:}, -c{:}, RZ;"
+STG  = "--:-:-:-:4      STG.E.SYS [clock_lo+{:}], e;"
+
+for i in range(loop_size):
+    SASS_CODE += [LDC.format(i)]
+    SASS_CODE += [CS2R.format(i+1)]
 
 for i in range(loop_size):
+    SASS_CODE += [IADD.format(i+1, i)]
     SASS_CODE += [STG.format(hex(i*4), i)]
 
 out_ = "\n" + "\n".join(SASS_CODE) + "\n"
diff --git a/sass_cubin/warp_schedule.sass b/sass_cubin/warp_schedule.sass
@@ -60,7 +60,7 @@ out_ = "\n" + "\n".join(SASS_CODE) + "\n"
 
 --:-:-:-:4      BAR.SYNC 0x0;
 
---:-:-:-:5     IADD3 e1, c2, -c1, RZ;
+--:-:-:-:5      IADD3 e1, c2, -c1, RZ;
 
 --:-:0:-:5      IMAD.WIDE clock_offset_lo, laneid, 0x4, clock_lo;
 --:-:-:-:4      MOV clock_offset_hi, clock_hi;
