I wonder if I just hit one of those circumstances.
Here is my story…
I have a dim3 grid. And I am clocking threadIdx.x == 0 of each block like this
namespace {
[...]
template<typename scalar_t>
__global__ void calculate_my_operation_kernel (
const scalar_t* __restrict__ x, const scalar_t* __restrict__ h,
scalar_t* __restrict__ output, const int method_idx,
long int* __restrict__ block_clock,
const int THREADS, const int C, const int PLANE_SIZE, const int PLANE_STEP,
const int INPUT_STEP, const int OUTPUT_STEP, const int CELL_STEP
) {
const int b = blockIdx.x;
const int f = blockIdx.y;
const int cell_idx = blockIdx.z * THREADS + threadIdx.x; // Position in PLANE
if(cell_idx >= PLANE_SIZE) return; // Ensures thread inside of PLANE_SIZE
// Indexes for memory addresses
const int x_idx_base = b*INPUT_STEP + cell_idx*CELL_STEP;
const int h_idx_base = f*INPUT_STEP + cell_idx*CELL_STEP;
// Here we create a dependency for calling the clock function before and
// after the operations we want to time. Having both of these ties is recommended.
// See https://forums.developer.nvidia.com/t/clock-function-on-titanx/44807/11
scalar_t out_; int c;
long int start = clock();
if (start != 0) out_ = 0.0; // Dependency on start
for (c = 0; c < C; c++) {
const int x_idx = x_idx_base + c*PLANE_STEP;
const scalar_t x_ = x[x_idx];
const int h_idx = h_idx_base + c*PLANE_STEP;
const scalar_t h_ = h[h_idx];
my_operation<scalar_t>[method_idx](x_, h_, &out_);
}
long int diff_;
if (c == C) diff_ = clock() - start; // Dependency on stop
const int output_idx = b*OUTPUT_STEP + f*PLANE_STEP + cell_idx*CELL_STEP;
output[output_idx] = out_;
// We just really care for the clock of one of the threads within a block.
const auto step_row = gridDim.z;
const auto step_plane = gridDim.y * gridDim.z;
if (threadIdx.x == 0) {
auto clock_idx = step_plane * blockIdx.x + step_row*blockIdx.y + blockIdx.z;
block_clock[clock_idx] = diff_;
}
}
} /** End of namespace */
Now, given the code above is from a torch CUDA kernel, and I have the data represented in torch::kLong as defined in the host function:
[...]
const auto Z = (H*W + THREADS - 1)/THREADS;
const dim3 GRID_SIZE(B, F, Z);
auto block_clock = torch::zeros({B, F, Z},
torch::TensorOptions().dtype(torch::kLong).device(
x.device().type(), x.device().index()
)
);
[...]
When visualizing the clock per block, I first get all clocks as a single numpy array:
# Clock shape is [128, 128, 16] for the picture below.
data = clock.cpu().numpy().view(np.uint64).reshape(-1,)
Then when visualizing it as a box plot, I have a max clock that is extremely outstanding, several orders of magnitude difference.
So, would this discrepancy in clock be due to one of those circumstances? Or I am casting the unsigned int values improperly somewhere?
Thank you for your time.
