support basic_shape on device

[lgarithm]

Currently the cuda grid system only supports up to 6d grid,

blockIdx.x
blockIdx.y
blockIdx.z
threadIdx.x
threadIdx.y
threadIdx.z

and it is tied to the kernel launch parameter.

But sometimes we would like to write something like

template <typename R, typename grid>
__global__ void k(grid g, const R* x, R* y)
{
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    coordinate c = coord(g, idx); // 
   // g(x,y,c);
}

template <typename R>
void f(const ttl::cuda_tensor_view<R> &x, const ttl::cuda_tensor_ref<R> &y){
    grid g = y.shape();
    constexpr int blocksPerGrid = 10;
    constexpr int threadsPerBlock = 10;
    k<R><<<blocksPerGrid, threadsPerBlock>>>(g, x, y);
}

[sjdrc]

Is this task for implementing ttl/algorithms in cuda kernels?

[lgarithm]

No, this is not related to the ttl/algorithm header.

ttl/algorithm will be some trivial algorithms for tensors. For non-trivial operators, they will be handled by stdnn-ops. I'm also developing the CUDA versions of those operators, but it's not open-sourced currently.

If you are looking for ttl/algorithms that works for ttl::cuda_tensor, you can use thrust. ttl/cuda_tensor should work well with thrust via the adaptors in #40.

[lgarithm]

template <std::size_t r> __device__ int get_coords(const int *g, int *c)
{
    const int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int j = idx;
    for (int i = r - 1; i >= 0; --i) {
        c[i] = j % g[i];
        j /= g[i];
    }
    return idx;
}