[QST] How to use an Iterator over parameters

[IzanCatalan]

What is your question?
Hi, I am trying to use Iterators to access filter parameters in Conv2dFprop.

In more detail, I am trying to access these params in this particular place:

cutlass/include/cutlass/conv/kernel/implicit_gemm_convolution.h

Line 332 in 24f991e

I try to use Iterator_B, which I believe iterates over the convolution filters.

The code I am using is the following:

while (iterator_B.valid()) {        
 typename Mma::IteratorB::Element value = *(iterator_B.get());
 printf("%f\n", static_cast<float>(value)); 
 iterator_B.advance();                    
}

But, I get some errors about that methons valid() and get() are not available in Conv2dFpropFilterTileAccessIteratorOptimized:

/mnt/beegfs/gap/[email protected]/cutlass/include/cutlass/conv/kernel/implicit_gemm_convolution.h(388): error: class "cutlass::conv::threadblock::TileIterator<cutlass::conv::threadblock::Conv2dFpropFilterTileAccessIteratorOptimized<cutlass::MatrixShape<32, 128>, ElementInputB, LayoutInputB, cutlass::transform::PitchLinearWarpRakedThreadMap<cutlass::PitchLinearShape<32, 128>, 128, cutlass::PitchLinearShape<4, 8>, 8>, cutlass::AlignedArray<ElementInputA, 8, 16>, false>>" has no member "valid"
/mnt/beegfs/gap/[email protected]/cutlass/include/cutlass/conv/kernel/implicit_gemm_convolution.h(389): error: class "cutlass::conv::threadblock::TileIterator<cutlass::conv::threadblock::Conv2dFpropFilterTileAccessIteratorOptimized<cutlass::MatrixShape<32, 128>, ElementInputB, LayoutInputB, cutlass::transform::PitchLinearWarpRakedThreadMap<cutlass::PitchLinearShape<32, 128>, 128, cutlass::PitchLinearShape<4, 8>, 8>, cutlass::AlignedArray<ElementInputA, 8, 16>, false>>" has no member "get"
2 errors detected in the compilation of "/mnt/beegfs/gap/[email protected]/cutlass/examples/16_ampere_tensorop_conv2dfprop/ampere_tensorop_conv2dfprop.cu".

How can I properly use IteratorB or access the pointer where the convolution filters are stored and access/modify them?

Thank you.

[hwu36]

I think it is actually easier for you to do it in a lower level - threadblock level:

https://github.com/NVIDIA/cutlass/blob/main/include/cutlass/conv/threadblock/implicit_gemm_multistage.h

it has many usages of iterator_B for you to use examples.

What you want to do seems can be inserted into the beginning of the epilogue. You need to make sure first without your change, the kernel can compile. If it cannot compile, it means this kernel is not supported.

[IzanCatalan]

Hi @hwu36 , I guess you mean, for example, to use the code near line 321, maybe it is working. Still, eventually, I discovered how you could access an modify the filter parameters without recurring to the iterator.
I modify the code and I have added in line332 of implicit_gemm_convolution.h the following:

        for (int n = 0; n < params.problem_size.K; n++) {
            for (int h = 0; h < params.problem_size.R; h++) {
                for (int w = 0; w < params.problem_size.S; w++) {
                    for (int c = 0; c < params.problem_size.C; c++) {
                        int index = 
                            n * (params.problem_size.R * params.problem_size.S * params.problem_size.C) +
                            h * (params.problem_size.S * params.problem_size.C) +
                            w * params.problem_size.C +
                            c;
                        params.ptr_B[index] = 2;
                        printf("B[%d, %d, %d, %d] = %f\n", n, h, w, c, static_cast<float>(params.ptr_B[index]));
                    }
                }
            }
        }

I deduce the shape of the filters by observing the several readme documents in the repo. Perhaps this method is not as technical as using iterators. Still, it allows me to parallelise the access to the filters in operator() function with several threads and to modify its values.

In any case, I would appreciate it if you could give me some insight about it.

[hwu36]

Yours is fine