How to get the execution time for the given function

[hulihan-start]

Hi! Recently I followed your docs and write a custom function, but I don't know how to evaluate the execution time.
Is there any document to tell me how to get the execution time of the given function? I saw the FreeTensor/test directory but found those are not what I want.

[roastduck]

The most simple way is to just use time.time from Python. Just run the function you are going to test for enough times, and the averaged time should be stable enough in most cases.

If you are look for a more accurate measurement, suppose f is the function compiled by FreeTensor, instead of calling f(some_args), you can do the following:

1. Call f.set_args(some_args).
2. Call average_time, standard_deviation = f.time(timing_rounds, warmup_rounds). Both arguments have default values so they can be omitted. The returned times are in milliseconds. This function is currently not documented at Python side, but it is actually a wrapper of this time function at C++ side.
3. Call return_values = f.collect_returns(). This step is required or there will be memory leak.

[roastduck]

By the way, if you are using the Python time.time() function to time a GPU program, you need to call device.sync() before time.time() to synchronize GPU streams, where device is a freetensor.Device object returned from something like freetensor.GPU().

[hulihan-start]

When I called f.collect_returns(), it returns a freetensor array as <freetensor_ffi.Array object at 0x7f75770767b0>.
And the function f.time() returns error as:
freetensor_ffi.DriverError: an illegal memory access was encountered
[WARNING] /data/not_backed_up/lihhu/FreeTensor_experiments/FreeTensor/include/driver.h:68: Return values must be collected, or there will be memory leaks

[hulihan-start]

I used time.time() before but I think there's something wrong with so long time execution time printed, so I want to know if you have any other ideas about the printing time.

[hulihan-start]

I found the problem, since I used numpy array before so the function may execute in CPU, when I changed to Tensor and mapped data on GPU, the execution time seems reasonable.

[roastduck]

When I called f.collect_returns(), it returns a freetensor array as <freetensor_ffi.Array object at 0x7f75770767b0>. And the function f.time() returns error as: freetensor_ffi.DriverError: an illegal memory access was encountered [WARNING] /data/not_backed_up/lihhu/FreeTensor_experiments/FreeTensor/include/driver.h:68: Return values must be collected, or there will be memory leaks

This seems like a real illegal memory access in the program. Please make sure your program can safely run for multiple times, especially when there are "inout" parameters.

I found the problem, since I used numpy array before so the function may execute in CPU, when I changed to Tensor and mapped data on GPU, the execution time seems reasonable.

The device to run the function is decided by Device used for compiling (for example by with ft.GPU()). As long as you use a GPU Device to compile, the program you run on a GPU. The program with directly using NumPy arrays as inputs is the data will be copied from CPU to GPU for every run. If you use a torch.Tensor that resides on a GPU, there will be no cross-device copy. If you don't want to use PyTorch, you can also use freetensor.Array, which copies the data lazily: cross-device copy only happens at the first run.

[hulihan-start]

Thanks a lot!

When I followed by the docs (https://roastduck.github.io/FreeTensor/guide/ad/) for Automatic Differentiation, an error was found:
TypeError: grad() got an unexpected keyword argument 'attach_backward'

After I removed the keyword attach_backward, another error was found:
TypeError: grad() missing 1 required positional argument: 'func'

I don't know if the docs need to update or if is there something wrong with my compilation.

[hulihan-start]

Another question is the shared memory usage.

I used to build a custom function that has matmul on GPU, but the generated code shows it doesn't use shared memory.

My custom function is :

    # @ft.grad(requires=['eemb', 'remb', 'proj'], provides=[ft.Return()])
    def transr(head: ft.Var[(batch_size,), "int64"], tail: ft.Var[(batch_size,), "int64"], relation: ft.Var[(batch_size,), "int64"], eemb: ft.Var[(enode, dim), "float32"], remb: ft.Var[(rnode, dim), "float32"], proj: ft.Var[(rnode, dim, dim), "float32"]):
        y = ft.empty((batch_size, dim), "float32")
        z = ft.empty((batch_size,), "float32")
        x = ft.empty((batch_size,), "float32")
        # for i in range(batch_size):
        #     for j in range(dim):
        #         for k in range(dim):
        #             y[i][j] += (eemb[head[i]][k] - eemb[tail[i]][k]) * proj[relation[i]][k][j]
        #         z[i] += ft.abs(y[i][j] + remb[relation[i]][j])
            # z[i] = x[i]
        for i in range(batch_size):
            y[i] = ft.libop.matmul((eemb[head[i]] - eemb[tail[i]]), proj[relation[i]]) + remb[relation[i]]
            for j in range(dim):
                x[i] += ft.abs(y[i][j])
            z[i] = x[i]
        return z

The generated code is below:

 1 
 2 #include <gpu_runtime.h>
 3 
 4 extern __shared__ uint8_t __shmem[];
 5 
 6 extern "C" {
 7 __global__ void __launch_bounds__(32 * 2 * 1) kernel0(mdspan_r<float, extents<4096>> _z, mdspan_r<const float, extents<51, 512>> _remb, mdspan_r<const int64_t, extents<4096>> _relation, mdspan_r<const float, extents<51, 512, 512>> _proj, mdspan_r<const int64_t, extents<4096>> _tail, mdspan_r<const int64_t, extents<4096>> _head, mdspan_r<const float, extents<93773, 512>> _eemb, __ByValArray<void *, 6> _params, uint8_t *__glmem) {
 8   auto _x = mdspan_r<float, extents<4096>>((float*)(__glmem + (0)));
 9   {
10     auto _out = mdspan_r<float, extents<32, 64, 2, 512>>((float*)(__glmem + (16384)));
11     for (int _i__1 = 0; _i__1 < 512; _i__1++) {
12       _out((int)threadIdx.x, (int)blockIdx.x, (int)threadIdx.y, _i__1) = (_eemb(_head((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))), _i__1) - _eemb(_tail((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))), _i__1));
13     }
14     {
15       auto _Y = mdspan_r<float, extents<32, 64, 2, 512>>((float*)(__glmem + (8404992)));
16       for (int _i__2 = 0; _i__2 < 512; _i__2++) {
17         _Y((int)threadIdx.x, (int)blockIdx.x, (int)threadIdx.y, _i__2) = 0;
18         for (int _i__3 = 0; _i__3 < 512; _i__3++) {
19           _Y((int)threadIdx.x, (int)blockIdx.x, (int)threadIdx.y, _i__2) += (_out((int)threadIdx.x, (int)blockIdx.x, (int)threadIdx.y, _i__3) * _proj(_relation((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))), _i__3, _i__2));
20         }
21       }
22       for (int _j = 0; _j < 512; _j++) {
23         _x((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))) += runtime_abs((_Y((int)threadIdx.x, (int)blockIdx.x, (int)threadIdx.y, _j) + _remb(_relation((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))), _j)));
24       }
25     }
26   }
27   _z((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y))) = _x((((64 * (int)blockIdx.x) + (int)threadIdx.x) + (32 * (int)threadIdx.y)));
28 }
29 
30 void run(void **__params, void **_returns, size_t **_retShapes, size_t *_retDims, GPUContext_t _ctx) {
31 __ByValArray<void *, 6> _params;
32 _params[0] = __params[0];
33 _params[1] = __params[1];
34 _params[2] = __params[2];
35 _params[3] = __params[3];
36 _params[4] = __params[4];
37 _params[5] = __params[5];
38 
39 cudaStream_t __stream = 0;
40 
41 uint8_t *__glmem = (uint8_t*)cudaNew(16793600);
42 
43 {
44   auto _head = mdspan_r<const int64_t, extents<4096>>((const int64_t*)(_params[0]));
45   auto _tail = mdspan_r<const int64_t, extents<4096>>((const int64_t*)(_params[1]));
46   auto _relation = mdspan_r<const int64_t, extents<4096>>((const int64_t*)(_params[2]));
47   auto _eemb = mdspan_r<const float, extents<93773, 512>>((const float*)(_params[3]));
48   auto _remb = mdspan_r<const float, extents<51, 512>>((const float*)(_params[4]));
49   auto _proj = mdspan_r<const float, extents<51, 512, 512>>((const float*)(_params[5]));
50   if (_returns[0] == NULL) {
51     _retShapes[0] = 1 > 0 ? (size_t*)malloc((_retDims[0] = 1) * sizeof(size_t)) : NULL;
52     checkCudaError(cudaMalloc(&_returns[0], (_retShapes[0][0] = 4096) * sizeof(float)));
53   }
54   auto _z = mdspan_r<float, extents<4096>>((float*)(_returns[0]));
55   checkCudaError(cudaFuncSetAttribute(kernel0, cudaFuncAttributeMaxDynamicSharedMemorySize, 0));
56   kernel0<<<dim3(64, 1, 1), dim3(32, 2, 1), 0, __stream>>>(_z, _remb, _relation, _proj, _tail, _head, _eemb, _params, __glmem);
57 }
58 
59 cudaFree(__glmem);
60 }
61 
62 }

Is there any tutorial that can tell me to use shared memory to generate high performance code ?

[roastduck]

Thanks a lot!

When I followed by the docs (https://roastduck.github.io/FreeTensor/guide/ad/) for Automatic Differentiation, an error was found: TypeError: grad() got an unexpected keyword argument 'attach_backward'

After I removed the keyword attach_backward, another error was found: TypeError: grad() missing 1 required positional argument: 'func'

I don't know if the docs need to update or if is there something wrong with my compilation.

The examples in docs is always tested in CI for every commit, so I don't think there will be a problem. Maybe you are running a out-of-date version of FreeTensor, and please pull a latest version from master. Sorry that we still haven't come out with a stable release.

[roastduck]

I used to build a custom function that has matmul on GPU, but the generated code shows it doesn't use shared memory.

I see your code has already in parallel. Did you use auto_schedule to parallelize it? auto_schedule is still in early development phase, and its shared memory support is very preliminary. To use shared memory in your program, it need to tile the variables fist according to GPU's shared memory size and other hardware parameters, and set each tile to use shared memory, which is currently not supported by auto_schedule.

To use shared memory, you need to modify the program by yourself, either by modify the code directly, or use schedules.

To specify shared memory in the code, you need to first tile the loops according to the GPU's hardware parameters. For variables in each tile, add a third argument to ft.empty. Set it to "gpu/shared" will set it to use GPU's shared memory. This is the most straight-forward way, but it makes the program no longer cross-platform.

You may also use schedules. You need to first use the split and reorder schedule to tile the loops. For each tile, use the cache schedule to create a sub-variable for this tile, and set it to use "gpu/shared". Unfortunately, without auto_schedule, the schedules may be complex, you can find https://github.com/roastduck/FreeTensor/blob/master/test/70.program/test_gpu_conv2d.py as an example of 2d-convolution with manual schedules. By the way, if you only need to set a variable to use "gpu/shared" without tiling, the set_mem_type schedule is enough. A list of all the schedules can be found here.