slower than cudnn

[Haoran-Solar]

cudnn int8_NHWC:

cutlass int4_NHWC:

seriously, cutlass int4 is faster than cudnn int8 on smaller size, but the performance is worse on larger ones.
how can i actually use cutlass in projects.

[hwu36]

Thank you for giving this a try.

There is a known performance bug in the epilogue. https://github.com/NVIDIA/cutlass/blob/master/include/cutlass/epilogue/threadblock/predicated_tile_iterator.h#L329-L333 needs to be changed to

        cutlass::arch::global_store<AccessType, sizeof(AccessType)>(
            frag_ptr[frag_row_idx * ThreadMap::Iterations::kColumn +
                     column],
            (void *)&memory_pointer[column * ThreadMap::Delta::kColumn /
                                    kElementsPerAccess],
            guard);

to use inline ptx to enforce vectorized store.

Also, would you please let us know more about your use case? Which hardware, how do you select cutlass kernels, what problem sizes, what frequencies/power limit, etc.? There are several ways to tune the performance depending on your use cases.

[Haoran-Solar]

I used the cases that you used in your unit test.
My hardware consists of T4 GPU card and Xeon Gold 6140 CPU. And no power limit is set.

[hwu36]

Thank you. Your information is very helpful. There is no good way to find one optimum kernel configuration that works well for all problem sizes. There are several things that can be tuned.

On Turing, the first thing you can tune is the threadblock/warp tile size (https://github.com/NVIDIA/cutlass/blob/master/examples/09_turing_tensorop_conv2dfprop/turing_tensorop_conv2dfprop.cu#L159-L162). The plausible tile sizes are here https://github.com/NVIDIA/cutlass/blob/master/tools/library/scripts/generator.py#L971-L976. The first tuple is threadblock tile sizes, the second is the stage number which is fixed to 2 for Turing, the third is warp count in each dimension. If threadblock size is 256x128x128, warp count is 4x2x1, the warp shape is then 64(=256/4)x64(=128/2)x128(=128/1).

The 2nd thing you can turn is the threadblock swizzling algorithms (

cutlass/examples/08_turing_tensorop_gemm/turing_tensorop_gemm.cu

Line 158 in ccb697b

using SwizzleThreadBlock = cutlass::gemm::threadblock::GemmIdentityThreadblockSwizzle<>; // <- ??

). You can try GemmIdentityThreadblockSwizzle<8>, GemmIdentityThreadblockSwizzle<4>, GemmIdentityThreadblockSwizzle<2>, GemmIdentityThreadblockSwizzle<1>, GemmIdentityThreadblockSwizzle.

You can use cutlass profiler to sweep different settings (https://github.com/NVIDIA/cutlass/blob/master/media/docs/profiler.md). Let us know if you need help on this.

Also, T4 has very low power package and all kernels are supposed to hit power throttling if not locking the frequency. To get reproducible result, we usually lock to frequency that is low enough that the kernels won't hit power throttling and run multiple iterations after some warmup run. Some useful commands are

nvidia-smi -i 0 -pm 1 // set to persistent mode
sudo nvidia-smi -lgc 900 -i 0 // lock the frequency to 900 MHz
nvidia-smi --format=csv --query-gpu=clocks.sm,temperature.gpu,fan.speed,power.draw,power.limit -lms 100 -i 0 // monitor gpu clock, temperature, etc in 100ms intervals. This needs to run in background during profiling to monitor the frequency is still locked.

Also, when you said "I used the cases that you used in your unit test.", Do you mean you run with the problem sizes in this file https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_problems.h ? Would you please let us know how do you run these in cudnn? What is even better is that you could use nvprof/nsight to let us know which kernel does cudnn run in your setting?

[Haoran-Solar]

Yes, I run with the problem sizes in file: https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_problems.h, resnet50 with 256 batch size.
About cudnn, I was comparing the efficiency between cudnn and cutlass. so I wrote a brief code.
Using CUDNN_TENSOR_NHWC and CUDNN_DATA_INT8 for input tensor, filter and output tensor. (cudnnSetTensor4dDescriptor(), sorry for some reason that is hard for me to download my code. )






Here is my test code.
And with nvprof, I found some interesting.

I might fail to use tensor core with int8. =_= But it's faster !?
By the way, following is my TimerGPU class.

[Haoran-Solar]

Is layout NCxHWx<32> same with CUDNN_TENSOR_NCHW_VECT_C?

[hwu36]

Is layout NCxHWx<32> same with CUDNN_TENSOR_NCHW_VECT_C?

Yes, if you use it together with CUDNN_DATA_INT8x32. See https://docs.nvidia.com/deeplearning/cudnn/developer-guide/index.html#nc32hw32-layout-x32 and https://docs.nvidia.com/deeplearning/cudnn/api/index.html#cudnnDataType_t

[manishucsd]

Hi @Haoran-Solar ,

Can you please run ResNet layers in cutlass profiler. Table 1 lists the top-performing kernel and layer parameter so you can copy it and run it with cutlass_profiler.

Steps to run int4 nhwc convolution with cutlass_profiler:

1. Get the source code:

git clone git@github.com:NVIDIA/cutlass.git
mkdir build; cd build

2. Build only int4 NHWC conv2d fprop kernels:

build$ cmake .. -DCUTLASS_NVCC_ARCHS='75' -DCUTLASS_LIBRARY_KERNELS="cutlass_tensorop*fprop*optimized*s4"

3. Lock the clocks:

nvidia-smi -i 0 -pm 1
sudo nvidia-smi -lgc 900 -i 0

4. Run a ResNet50 layer 17 from Table 1.

build$ ./tools/profiler/cutlass_profiler --kernels=cutlass_tensorop_s4_i8832fprop_optimized_s4_256x128_128x2_nhwc --n=256 --h=7 --w=7 --c=512 --k=512 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1

Can you measure the runtimes using the steps above? You should be able to put information in Table 1 in a script and run all the layers at once to generate the .csv using cutlass_profiler options --output=int4_cutlass_perf.csv --append=true

Tip: Just run --kernels=cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc for all resnet layer as 128x128_128x2_nhwc wins for most layers. I did a sweep over all tile sizes to find the winning tile size for a layer as reported in Table 1.

ResnetLayer	TopPerformingKernel	BatchSize	LayerParams
1	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=64 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
2	cutlass_tensorop_s4_i8832fprop_optimized_s4_64x64_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=64 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
3	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=64 --k=64 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
4	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=256 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
5	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=256 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
6	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=56 --w=56 --c=256 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
7	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=28 --w=28 --c=128 --k=128 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
8	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=28 --w=28 --c=128 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
9	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=28 --w=28 --c=512 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
10	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=28 --w=28 --c=512 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
11	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=28 --w=28 --c=512 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
12	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=14 --w=14 --c=256 --k=256 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
13	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=14 --w=14 --c=256 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
14	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=14 --w=14 --c=1024 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
15	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=14 --w=14 --c=1024 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
16	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=14 --w=14 --c=1024 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1
17	cutlass_tensorop_s4_i8832fprop_optimized_s4_256x128_128x2_nhwc	256	--n=256 --h=7 --w=7 --c=512 --k=512 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
18	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=7 --w=7 --c=512 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1
19	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	256	--n=256 --h=7 --w=7 --c=2048 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1

Table 1.

[JerryBay]

I see two different IteratorAlgorithm. kAnalytic and kOptimized. You say kOptimized is faster. But is any extra cost on data rearrangment like int8x32 in cudnn.

[hwu36]

I see two different IteratorAlgorithm. kAnalytic and kOptimized. You say kOptimized is faster. But is any extra cost on data rearrangment like int8x32 in cudnn.

IteratorAlgorithm is orthogonal to the input data layout.

https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_fprop_implicit_gemm_s8nhwc_s8nhwc_s32nhwc_tensor_op_s32_sm75.cu and https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_fprop_implicit_gemm_s4nhwc_s4nhwc_s32nhwc_tensor_op_s32_sm75.cu have the Analytic and Optimized versions for NHWC.

https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_fprop_implicit_gemm_s8ncxhwx_s8cxrskx_s8ncxhwx_tensor_op_s32_sm75.cu and https://github.com/NVIDIA/cutlass/blob/master/test/unit/conv/device/conv2d_fprop_implicit_gemm_s4ncxhwx_s4cxrskx_s4ncxhwx_tensor_op_s32_sm75.cu have the Analytic and Optimized versions for NC3xHWx (x=32 for int8 and x=64 for int4). NCxHWx needs to reorder the filter channels if the outputs are int8 or int4 to bypass the shared memory in the epilogue.

[manishucsd]

I see two different IteratorAlgorithm. kAnalytic and kOptimized. You say kOptimized is faster. But is any extra cost on data rearrangment like int8x32 in cudnn.

The following subsection of ImplicitGemmConvolution documentation explains analytic and optimized iterators: https://github.com/NVIDIA/cutlass/blob/master/media/docs/implicit_gemm_convolution.md#cutlass-components

[manishucsd]

Hi @Haoran-Solar,

Please find the int4 (S4) runtimes on device Tesla T4 - SM 7.5, 40 SMs (you can print similar device details ./tools/profiler/cutlass_profiler --help OR nvidia-smi).

The numbers I observe with the latest 11.2 CUDA toolkit are much better than what you reported in comment 1.

Let us know if you are able to repro the below CUTLASS runtimes for (S4 <= S4*S4 + S32 NHWC) using steps in #154 (comment) ?

DataType	chip	conv_kind	LayerID	TopOperation	LayerParams	Runtime (ms)
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	1	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	1.90826
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	2	cutlass_tensorop_s4_i8832fprop_optimized_s4_64x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.592426
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	3	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=64 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	1.7911
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	4	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.866032
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	5	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	1.27112
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	6	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.331363
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	7	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=128 --k=128 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.691952
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	8	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=128 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.96567
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	9	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.466544
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	10	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.892182
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	11	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.25161
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	12	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=256 --k=256 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.644016
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	13	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=256 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.639203
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	14	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.395565
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	15	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.699488
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	16	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.18095
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	17	cutlass_tensorop_s4_i8832fprop_optimized_s4_256x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=512 --k=512 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.57552
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	18	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=512 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.452774
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	19	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=2048 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.302957

[Haoran-Solar]

Thanks for your generous help. I don't have root so no frequency limit can be set. But I changed the iterator algorithm to optimized and GemmIdentityThreadblockSwizzle<4> to have better performance.
I will try cutlass_profiler later and let you know about runtimes.
And although I changed kernel to have better performance in most layer , but layer 1 and 2 still have bad performance.

[hwu36]

@Haoran-Solar ,

First, have you tried to use different tile sizes? For example, @manishucsd 's data uses threadblock_64x64_warp_32x32 tile size in layer two. It is normal that every layer uses different kernels.

Second, int4 has worse quantization problem than int8. On Turing, int4 tensor core kernels process 128 elements in the RSC dimension in one main loop iteration and int8 tensor core kernels process 64 elements in the RSC dimension in one main loop iteration. Every loop iteration takes the same amount of time in theory for int4 and int8 kernels. In layer one and layer two, RSC = 64 which needs only one iteration for both int4 and int8 kernels, so that int4 kernels have no advantage in these two layers. What is your end-to-end speedup?

Third, as you noticed that cuDNN does not choose to use tensor core kernels for whatever reasons, you may want to consider using cutlass int8 tensor core kernels.

[Haoran-Solar]

Hi @manishucsd ,
Here is my runtimes.

DataType	chip	conv_kind	LayerID	Kernel	LayerParams	Runtimes(ms)
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	1	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	1.39524
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	2	cutlass_tensorop_s4_i8832fprop_optimized_s4_64x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.376495
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	3	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=64 --k=64 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	1.29968
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	4	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x64_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=64 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.738804
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	5	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.956397
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	6	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=56 --w=56 --c=256 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.215717
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	7	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=128 --k=128 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.48018
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	8	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=128 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.663186
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	9	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=128 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.318629
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	10	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.667512
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	11	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=28 --w=28 --c=512 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.163109
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	12	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=256 --k=256 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.431603
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	13	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=256 --k=1024 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.418936
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	14	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=256 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.263654
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	15	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.503855
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	16	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=14 --w=14 --c=1024 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=2 --stride_w=2 --dilation_h=1 --dilation_w=1	0.115354
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	17	cutlass_tensorop_s4_i8832fprop_optimized_s4_256x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=512 --k=512 --r=3 --s=3 --pad_h=1 --pad_w=1 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.367956
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	18	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=512 --k=2048 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.288225
S4 <= S4*S4 + S32 (NHWC)	turing	fprop	19	cutlass_tensorop_s4_i8832fprop_optimized_s4_128x128_128x2_nhwc	--n=256 --h=7 --w=7 --c=2048 --k=512 --r=1 --s=1 --pad_h=0 --pad_w=0 --stride_h=1 --stride_w=1 --dilation_h=1 --dilation_w=1	0.195773

I didn't set the frequency limit(no root), and I got better performance with 10.2 CUDA toolkit.

[hwu36]

Thank you, @Haoran-Solar . Your latest int4 performance number w/o locking frequency is much better than any other int4 or int8 performance number reported in this thread. I guess you are unblocked now. We are happy to help if you have other questions when using CUTLASS.

P.S. CUDA 10.2 is indeed a very good compiler for Turing.