0
Я действительно не знаю, как назвать проблему, с которой я столкнулся, поэтому любой мод переименует ее соответственно, если вы так считаете.Проблема производительности CUDA
я получил следующую матрицу векторного умножения ядра:
__global__ void dmv_gpu_shmem(const value_t *a, const value_t *x, value_t *y,
size_t n)
{
extern __shared__ value_t shmem_buf[];
int ltid = threadIdx.x;
int gtid = get_global_tid();
value_t _y = 0.0;
if (gtid > n)
return;
int last_id = n/blockDim.x;
for(size_t j=0; j< last_id; j++) {
shmem_buf[ltid] = x[blockDim.x*j + ltid];
__syncthreads();
for(size_t i=0; i< blockDim.x; i++) {
_y += a[gtid + (i + j*blockDim.x)*n] * shmem_buf[i];
}
__syncthreads();
}
y[gtid] = _y;
}
Я должен проверить это ядро в течение некоторого размера блока, но результаты синхронизации я получил где не так хорошо. Поэтому я решил изменить это следующим образом и зафиксировать размер блока 32 при вызове ядра
__global__ void dmv_gpu_shmem(const value_t *a, const value_t *x, value_t *y,
size_t n)
{
extern __shared__ value_t shmem_buf[];
int ltid = threadIdx.x;
int gtid = get_global_tid();
value_t _y = 0.0;
if (gtid > n)
return;
int last_id = n/32;
for(size_t j=0; j< last_id; j++) {
shmem_buf[ltid] = x[32*j + ltid];
__syncthreads();
for(size_t i=0; i< 32; i++) {
_y += a[gtid + (i + j*32)*n] * shmem_buf[i];
}
__syncthreads();
}
y[gtid] = _y;
}
К моему большому удивлению, ядро не получил несколько раз лучше с точки зрения времени выполнения, и я совершенно не знаю, почему это происходит.
Может ли кто-нибудь более опытный объяснить это?
Также, учитывая эту ситуацию, как я могу получить максимум из моего ядра, используя все различные блоки, которые я хочу? Я не могу этого сделать для всех ...
размера блокаEdit:
Это должен быть рабочим репро случай:
#include <stdlib.h>
#include <stdio.h>
#include <sys/time.h>
#include <cuda.h>
#include "cublas_v2.h" //CUBLAS LIBRARY
#ifndef VALUES_MAX
# define VALUES_MAX 1.
#endif
#ifndef EPS
# define EPS 1.e-6
#endif
#ifndef NR_ITER
# define NR_ITER 200
#endif
enum
{
GPU_NAIVE = 0,
GPU_COALESCED,
GPU_SHMEM,
GPU_KERNEL_END
};
void *gpu_alloc(size_t count)
{
void *ret;
if (cudaMalloc(&ret, count) != cudaSuccess) {
ret = NULL;
}
return ret;
}
int copy_to_gpu(const void *host, void *gpu, size_t count)
{
if (cudaMemcpy(gpu, host, count, cudaMemcpyHostToDevice) != cudaSuccess)
return -1;
return 0;
}
int copy_from_gpu(void *host, const void *gpu, size_t count)
{
if (cudaMemcpy(host, gpu, count, cudaMemcpyDeviceToHost) != cudaSuccess)
return -1;
return 0;
}
void mat_init_rand(float **a, size_t n, float max)
{
size_t i, j;
for (i = 0; i < n; ++i)
{
for (j = 0; j < n; ++j)
{
//printf("%d %d \n", i, j);
a[i][j] = 2 * (((float) drand48()) - 0.5) * max;
}
}
}
void vec_init(float *v, size_t n, float val)
{
size_t i;
for (i = 0; i < n; ++i)
{
v[i] = val;
}
}
void vec_init_rand(float *v, size_t n, float max)
{
size_t i;
for (i = 0; i < n; ++i)
{
v[i] = 2 * (((float) drand48()) - 0.5) * max;
}
}
void vec_print(const float *v, size_t n)
{
size_t i;
for (i = 0; i < n; ++i)
printf("%f \n", v[i]);
}
void **calloc_2d(size_t n, size_t m, size_t size)
{
char **ret = (char **) malloc(n*sizeof(char *));
if (ret) {
char *area = (char *) calloc(n*m, size);
if (area) {
for (size_t i = 0; i < n; ++i)
ret[i] = (char *) &area[i*m*size];
} else {
free(ret);
ret = NULL;
}
}
return (void **) ret;
}
void **copy_2d(void **dst, const void **src, size_t n, size_t m, size_t size)
{
memcpy(dst[0], src[0], n*m*size);
return dst;
}
void free_2d(void **array)
{
free(array[0]);
free(array);
}
__global__ void dmv_gpu_shmem(const float *a, const float *x, float *y,
size_t n)
{
extern __shared__ float shmem_buf[];
int ltid = threadIdx.x;
int gtid = blockIdx.x*blockDim.x+threadIdx.x;
float _y = 0.0;
if (gtid > n)
return;
int last_id = n/blockDim.x;
for(size_t j=0; j< last_id; j++) {
shmem_buf[ltid] = x[blockDim.x*j + ltid];
__syncthreads();
for(size_t i=0; i< blockDim.x; i++) {
_y += a[gtid + (i + j*blockDim.x)*n] * shmem_buf[i];
}
__syncthreads();
}
y[gtid] = _y;
}
__global__ void dmv_gpu_shmem_static(const float *a, const float *x, float *y,
size_t n)
{
extern __shared__ float shmem_buf[];
int ltid = threadIdx.x;
int gtid = blockIdx.x*blockDim.x+threadIdx.x;
float _y = 0.0;
if (gtid > n)
return;
int last_id = n/32;
for(size_t j=0; j< last_id; j++) {
shmem_buf[ltid] = x[32*j + ltid];
__syncthreads();
for(size_t i=0; i< 32; i++) {
_y += a[gtid + (i + j*32)*n] * shmem_buf[i];
}
__syncthreads();
}
y[gtid] = _y;
}
int main(int argc, char **argv)
{
if (argc < 2) {
printf("Wrong arguments \n");
return -1;
}
size_t orig_n = atoi(argv[1]);
/* Read block size and kernel to launch from the environment */
const char *env_gpu_kernel = getenv("GPU_KERNEL");
const char *env_gpu_block_size = getenv("GPU_BLOCK_SIZE");
int kernel = (env_gpu_kernel) ? atoi(env_gpu_kernel) : GPU_NAIVE;
int block_size = (env_gpu_block_size) ? atoi(env_gpu_block_size) : 256;
//Adjust Matrix to fit blocksize
size_t n = ((orig_n - 1)/block_size + 1)*block_size;
int grid_size = (n-1)/block_size + 1;
printf("Matrix size: %zd\n", orig_n);
printf("Input Block size: %zd\n", block_size);
printf("Adjusted matrix size: %zd\n", n);
/*
* Allocate the structures.
*
* Initialization to zero is crucial if you adjusted the matrix
* size.
*/
float **A = (float **) calloc_2d(n, n, sizeof(**A));
float *x = (float *) calloc(n, sizeof(*x));
float *y = (float *) calloc(n, sizeof(*y));
/* Initialize */
srand48(0);
mat_init_rand(A, orig_n, VALUES_MAX);
vec_init_rand(x, orig_n, VALUES_MAX);
vec_init(y, orig_n, 0.0);
printf("Setup Complete\n");
/*
* FILLME: Set up the blocks, grid and shared memory depending on
* the kernel. Make any transformations to the input
* matrix here.
*/
//Transposing Matrix for Shared and Coalesced Matrices
float tmp;
for(size_t i=0;i<n;i++)
for(size_t j=i+1;j<n;j++) {
tmp=A[i][j];
A[i][j] = A[j][i];
A[j][i] = tmp;
}
dim3 gpu_block(block_size, 1); // Number of threads
dim3 gpu_grid(grid_size, 1); // Number of blocks
size_t shmem_size = 0; // Shared memory size
/* Set SHARED MEMORY size */
shmem_size = block_size * sizeof(float);
printf(">>>> Begin of record <<<<\n");
printf("Block size: %dx%d\n", gpu_block.x, gpu_block.y);
printf("Grid size : %dx%d\n", gpu_grid.x, gpu_grid.y);
printf("Shared memory size: %ld bytes\n", shmem_size);
/* GPU allocations */
float *gpu_A = (float *) gpu_alloc(n*n*sizeof(*gpu_A));
float *gpu_x = (float *) gpu_alloc(n*sizeof(*gpu_x));
float *gpu_y = (float *) gpu_alloc(n*sizeof(*gpu_y));
/* Copy data to GPU */
copy_to_gpu(A[0], gpu_A, n*n*sizeof(*gpu_A));
copy_to_gpu(x, gpu_x, n*sizeof(*gpu_x));
/* Reset y and copy it to GPU */
vec_init(y, n, 0.0);
copy_to_gpu(y, gpu_y, n*sizeof(*gpu_y));
dmv_gpu_shmem<<<gpu_grid,gpu_block,shmem_size>>>
(gpu_A, gpu_x, gpu_y, n);
if (cudaGetLastError() != cudaSuccess)
printf("gpu kernel failed to launch \n");
dmv_gpu_shmem_static<<<gpu_grid,gpu_block,shmem_size>>>
(gpu_A, gpu_x, gpu_y, n);
if (cudaGetLastError() != cudaSuccess)
printf("gpu kernel failed to launch \n");
cudaDeviceSynchronize();
/* Free resources on host */
free_2d((void **) A);
free(x);
free(y);
/* Free resources on GPU */
cudaFree(gpu_A);
cudaFree(gpu_x);
cudaFree(gpu_y);
return EXIT_SUCCESS;
}
Compile с
nvcc dmv_test_case.cu
Выполнить с
GPU_KERNEL=2 GPU_BLOCK_SIZE=32 ./a.out 2048
G Переменная PU_KERNEL в этом случае ничего не делает. GPU_BLOCK_SIZE очевиден. Число аргумент размер вектора (п) и размер матрицы (пхп)
Возможно компилятор оптимизации из-за подсчета фиксированного цикла. Тем не менее, будет полезен полный пример воспроизведения. – talonmies
Не nvcc делает автоматическую оптимизацию на ядрах, я думаю, что прочитал это в руководстве по программированию. Я попытаюсь сделать репродуцирование, но это слишком сложно, так как вся программа снова разбита на 5 файлов ... –
Да, но петли могут быть развернуты во втором ядре, а не в первом из-за постоянного подсчета поездок , – talonmies