Основываясь на моем исследовании, есть 2 разные стратегии для реализации мозаичной версии свертки с cuda.Я хочу узнать больше об этом, и хотел бы увидеть, как они соотносятся друг с другом, каковы преимущества и недостатки каждой стратегии и как выбрать.Ниже приведены реализации двух разных стратегий.
Стратегия 1: размер плитки совпадает с размером выходного файла и требует несколько шагов для загрузки ввода.
#define MASK_WIDTH 3
#define MASK_RADIUS 1
#define TILE_WIDTH 8
#define SHAREDMEM_DIM (TILE_WIDTH + (MASK_RADIUS * 2))
__constant__ float deviceMask[MASK_WIDTH * MASK_WIDTH * MASK_WIDTH];
__global__ void conv3d(float *inputArray,
float *outputArray,
const int z_size,
const int y_size,
const int x_size) {
__shared__ float subTile[SHAREDMEM_DIM][SHAREDMEM_DIM][SHAREDMEM_DIM];
int bx = blockIdx.x, tx = threadIdx.x;
int by = blockIdx.y, ty = threadIdx.y;
int bz = blockIdx.z, tz = threadIdx.z;
int destination = (tz * TILE_WIDTH * TILE_WIDTH) + (ty * TILE_WIDTH) + tx;
int destTmp = destination;
int dX = destTmp % SHAREDMEM_DIM;
destTmp = destTmp / SHAREDMEM_DIM;
int dY = destTmp % SHAREDMEM_DIM;
destTmp = destTmp / SHAREDMEM_DIM;
int dZ = destTmp;
int inputZ = dZ + (bz * TILE_WIDTH) - MASK_RADIUS;
int inputY = dY + (by * TILE_WIDTH) - MASK_RADIUS;
int inputX = dX + (bx * TILE_WIDTH) - MASK_RADIUS;
int input = (inputZ * y_size * x_size) + (inputY * x_size) + inputX;
if( inputZ >= 0 && inputZ < z_size
&& inputY >= 0 && inputY < y_size
&& inputX >= 0 && inputX < x_size){
subTile[dZ][dY][dX] = inputArray[input];
}
else{
subTile[dZ][dY][dX] = 0;
}
destination = TILE_WIDTH * TILE_WIDTH * TILE_WIDTH
+ (tz * TILE_WIDTH * TILE_WIDTH) + (ty * TILE_WIDTH) + tx;
destTmp = destination;
dX = destTmp % SHAREDMEM_DIM;
destTmp = destTmp / SHAREDMEM_DIM;
dY = destTmp % SHAREDMEM_DIM;
destTmp = destTmp / SHAREDMEM_DIM;
dZ = destTmp;
inputZ = dZ + (bz * TILE_WIDTH) - MASK_RADIUS;
inputY = dY + (by * TILE_WIDTH) - MASK_RADIUS;
inputX = dX + (bx * TILE_WIDTH) - MASK_RADIUS;
input = (inputZ * y_size * x_size) + (inputY * x_size) + inputX;
if(dZ < SHAREDMEM_DIM){
if( inputZ >= 0 && inputZ < z_size
&& inputY >= 0 && inputY < y_size
&& inputX >= 0 && inputX < x_size ) {
subTile[dZ][dY][dX] = inputArray[input];
}
else{
subTile[dZ][dY][dX] = 0;
}
}
__syncthreads();
float sum = 0;
int z, y, x;
for(z = 0; z < MASK_WIDTH; z++){
for(y = 0; y < MASK_WIDTH; y++){
for(x = 0; x < MASK_WIDTH; x++){
sum += subTile[tz + z][ty + y][tx + x]
* deviceMask[x + (y * MASK_WIDTH) + (z * MASK_WIDTH * MASK_WIDTH)];
}
}
}
z = tz + (bz * TILE_WIDTH);
y = ty + (by * TILE_WIDTH);
x = tx + (bx * TILE_WIDTH);
if(z < z_size && y < y_size && x < x_size){
outputArray[x + (y * x_size) + (z * y_size * x_size)] = sum;
}
__syncthreads();
}
Вторая стратегияустановить одинаковый размер блока с плиткой ввода.При вычислении вывода некоторые потоки отключаются.
#define TILE_X 14
#define TILE_Y 6
#define TILE_Z 6
#define MASK_WIDTH 3
#define MASK_SIZE MASK_WIDTH * MASK_WIDTH * MASK_WIDTH
__constant__ float mask[MASK_WIDTH][MASK_WIDTH][MASK_WIDTH];
__global__ void conv3d(float *input, float *output, const int z_size, const int y_size, const int x_size) {
__shared__ float inputTile [TILE_Z+MASK_WIDTH-1][TILE_Y+MASK_WIDTH-1][TILE_X+MASK_WIDTH-1];
int tx = threadIdx.x; int ty = threadIdx.y; int tz = threadIdx.z;
int bx = blockIdx.x; int by = blockIdx.y; int bz = blockIdx.z;
int x_o = bx * TILE_X + tx
int y_o = by * TILE_Y + ty;
int z_o = bz * TILE_Z + tz;
int x_i = x_o - MASK_WIDTH/2;
int y_i = y_o - MASK_WIDTH/2;
int z_i = z_o - MASK_WIDTH/2;
if (x_i >= 0 && y_i >= 0 && z_i >= 0 && x_i < x_size && y_i < y_size && z_i < z_size)
inputTile[tz][ty][tx] = input[(z_i * y_size + y_i) * x_size + x_i];
else
inputTile[tz][ty][tx] = 0.0;
__syncthreads();
float acc = 0.0;
if(tz < TILE_Z && ty < TILE_Y && tx < TILE_X) {
for(int z_mask = 0; z_mask < Z_MASK_WIDTH; z_mask++) {
for(int y_mask = 0; y_mask < Y_MASK_WIDTH; y_mask++) {
for(int x_mask = 0; x_mask < X_MASK_WIDTH; x_mask++) {
acc += mask[z_mask][y_mask][x_mask] *
inputTile[tz+z_mask][ty+y_mask][tx+x_mask];
}
}
}
if(z_o < z_size && y_o < y_size && x_o < x_size)
output[(z_o * y_size + y_o) * x_size + x_o] = acc;
}
}
Есть идеи о том, как выбирать между ними?Кроме того, какая версия чаще всего используется на практике, например, при глубоком обучении?Также, если у вас есть какие-либо комментарии к коду, пожалуйста, дайте мне знать!