@anatoliylitv

I tried with assert in vectorized elementwise kernel by adding CUDA_KERNEL_ASSERT(bws * blockIdx.x <= N); here
Didn't see any perf impact (1-4% variation)

for 16x16: < 1%
for 4096x4096: < 4%

import time
import statistics
import gc
import torch

device = "cuda"
for SIZE in [16, 4096, 8192, 16384] :        #, 19737]:
    f_t = []
    for _ in range (5):
        a = torch.randn(SIZE, SIZE, device=device)
        b = torch.randn(SIZE, SIZE, device=device)

        for _ in range(20):
            torch.add(a, b)
        torch.cuda.synchronize()

        start = time.perf_counter_ns()
        for _ in range(100):
            torch.add(a, b)
        torch.cuda.synchronize()
        elapsed = time.perf_counter_ns() - start
        avg_t = elapsed / (1000 * 100)
        f_t.append(avg_t)
        #print(f"device: {device}")
        #print(f"total: {elapsed * 1000:.3f} ms over 100 iters")
        print(f"size: {SIZE}x{SIZE} , avg: {avg_t:.3f} us/iter")
        time.sleep(1)
        gc.collect()
        torch.cuda.empty_cache()
    print(f"size: {SIZE}x{SIZE} , tavg: {statistics.mean(f_t):.3f} us")
    print("======================================")

Gather information and generate reproducers about assert in kernels performance drop. #3151

Description

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions