libgpuarray/src/gpuarray_buffer_cuda.c at master · jrao1/libgpuarray · GitHub

History

1710 lines (1490 loc) · 45.8 KB

Raw

100

101

102

103

104

105

106

107

108

109

110

111

112

113

114

115

116

117

118

119

120

121

122

123

124

125

126

127

128

129

130

131

132

133

134

135

136

137

138

139

140

141

142

143

144

145

146

147

148

149

150

151

152

153

154

155

156

157

158

159

160

161

162

163

164

165

166

167

168

169

170

171

172

173

174

175

176

177

178

179

180

181

182

183

184

185

186

187

188

189

190

191

192

193

194

195

196

197

198

199

200

201

202

203

204

205

206

207

208

209

210

211

212

213

214

215

216

217

218

219

220

221

222

223

224

225

226

227

228

229

230

231

232

233

234

235

236

237

238

239

240

241

242

243

244

245

246

247

248

249

250

251

252

253

254

255

256

257

258

259

260

261

262

263

264

265

266

267

268

269

270

271

272

273

274

275

276

277

278

279

280

281

282

283

284

285

286

287

288

289

290

291

292

293

294

295

296

297

298

299

300

301

302

303

304

305

306

307

308

309

310

311

312

313

314

315

316

317

318

319

320

321

322

323

324

325

326

327

328

329

330

331

332

333

334

335

336

337

338

339

340

341

342

343

344

345

346

347

348

349

350

351

352

353

354

355

356

357

358

359

360

361

362

363

364

365

366

367

368

369

370

371

372

373

374

375

376

377

378

379

380

381

382

383

384

385

386

387

388

389

390

391

392

393

394

395

396

397

398

399

400

401

402

403

404

405

406

407

408

409

410

411

412

413

414

415

416

417

418

419

420

421

422

423

424

425

426

427

428

429

430

431

432

433

434

435

436

437

438

439

440

441

442

443

444

445

446

447

448

449

450

451

452

453

454

455

456

457

458

459

460

461

462

463

464

465

466

467

468

469

470

471

472

473

474

475

476

477

478

479

480

481

482

483

484

485

486

487

488

489

490

491

492

493

494

495

496

497

498

499

500

501

502

503

504

505

506

507

508

509

510

511

512

513

514

515

516

517

518

519

520

521

522

523

524

525

526

527

528

529

530

531

532

533

534

535

536

537

538

539

540

541

542

543

544

545

546

547

548

549

550

551

552

553

554

555

556

557

558

559

560

561

562

563

564

565

566

567

568

569

570

571

572

573

574

575

576

577

578

579

580

581

582

583

584

585

586

587

588

589

590

591

592

593

594

595

596

597

598

599

600

601

602

603

604

605

606

607

608

609

610

611

612

613

614

615

616

617

618

619

620

621

622

623

624

625

626

627

628

629

630

631

632

633

634

635

636

637

638

639

640

641

642

643

644

645

646

647

648

649

650

651

652

653

654

655

656

657

658

659

660

661

662

663

664

665

666

667

668

669

670

671

672

673

674

675

676

677

678

679

680

681

682

683

684

685

686

687

688

689

690

691

692

693

694

695

696

697

698

699

700

701

702

703

704

705

706

707

708

709

710

711

712

713

714

715

716

717

718

719

720

721

722

723

724

725

726

727

728

729

730

731

732

733

734

735

736

737

738

739

740

741

742

743

744

745

746

747

748

749

750

751

752

753

754

755

756

757

758

759

760

761

762

763

764

765

766

767

768

769

770

771

772

773

774

775

776

777

778

779

780

781

782

783

784

785

786

787

788

789

790

791

792

793

794

795

796

797

798

799

800

801

802

803

804

805

806

807

808

809

810

811

812

813

814

815

816

817

818

819

820

821

822

823

824

825

826

827

828

829

830

831

832

833

834

835

836

837

838

839

840

841

842

843

844

845

846

847

848

849

850

851

852

853

854

855

856

857

858

859

860

861

862

863

864

865

866

867

868

869

870

871

872

873

874

875

876

877

878

879

880

881

882

883

884

885

886

887

888

889

890

891

892

893

894

895

896

897

898

899

900

901

902

903

904

905

906

907

908

909

910

911

912

913

914

915

916

917

918

919

920

921

922

923

924

925

926

927

928

929

930

931

932

933

934

935

936

937

938

939

940

941

942

943

944

945

946

947

948

949

950

951

952

953

954

955

956

957

958

959

960

961

962

963

964

965

966

967

968

969

970

971

972

973

974

975

976

977

978

979

980

981

982

983

984

985

986

987

988

989

990

991

992

993

994

995

996

997

998

999

1000

#define _CRT_SECURE_NO_WARNINGS

#include "private.h"

#include "private_cuda.h"

#include "loaders/libnvrtc.h"

#include "loaders/libcublas.h"

#include <sys/types.h>

#include <assert.h>

#include <stdlib.h>

#include <cache.h>

#include "util/strb.h"

#include "util/xxhash.h"

#include "gpuarray/buffer.h"

#include "gpuarray/util.h"

#include "gpuarray/error.h"

#include "gpuarray/buffer_blas.h"

#include "gpuarray/extension.h"

STATIC_ASSERT(DONTFREE == GPUARRAY_CUDA_CTX_NOFREE, cuda_nofree_eq);

STATIC_ASSERT(CUDA_WAIT_READ == GPUARRAY_CUDA_WAIT_READ, cuda_wait_read_eq);

STATIC_ASSERT(CUDA_WAIT_WRITE == GPUARRAY_CUDA_WAIT_WRITE, cuda_wait_write_eq);

STATIC_ASSERT(sizeof(GpuArrayIpcMemHandle) == sizeof(CUipcMemHandle), cuda_ipcmem_eq);

/* Allocations will be made in blocks of at least this size */

#define BLOCK_SIZE (4 * 1024 * 1024)

/* No returned allocations will be smaller than this size. Also, they

* will be aligned to this size.

* Some libraries depend on this value and will crash if it's smaller.

#define FRAG_SIZE (64)

static CUresult err;

GPUARRAY_LOCAL const gpuarray_buffer_ops cuda_ops;

static void cuda_freekernel(gpukernel *);

static int cuda_property(gpucontext *, gpudata *, gpukernel *, int, void *);

static int cuda_waits(gpudata *, int, CUstream);

static int cuda_records(gpudata *, int, CUstream);

static int detect_arch(const char *prefix, char *ret, CUresult *err);

static gpudata *new_gpudata(cuda_context *ctx, CUdeviceptr ptr, size_t size);

static int strb_eq(void *_k1, void *_k2) {

strb *k1 = (strb *)_k1;

strb *k2 = (strb *)_k2;

return (k1->l == k2->l &&

memcmp(k1->s, k2->s, k1->l) == 0);

}

static uint32_t strb_hash(void *_k) {

strb *k = (strb *)_k;

return XXH32(k->s, k->l, 42);

}

static int setup_done = 0;

static int major = -1;

static int minor = -1;

static int setup_lib(void) {

int res, tmp;

const char *ver;

if (!setup_done) {

res = load_libcuda();

if (res != GA_NO_ERROR)

return res;

err = cuInit(0);

if (err != CUDA_SUCCESS)

return GA_IMPL_ERROR;

ver = getenv("GPUARRAY_CUDA_VERSION");

if (ver == NULL || strlen(ver) != 2) {

err = cuDriverGetVersion(&tmp);

if (err != CUDA_SUCCESS)

return GA_IMPL_ERROR;

major = tmp / 1000;

minor = (tmp / 10) % 10;

} else {

major = ver[0] - '0';

minor = ver[1] - '0';

}

if (major > 9 || major < 0 || minor > 9 || minor < 0)

return GA_VALUE_ERROR;

res = load_libnvrtc(major, minor);

if (res != GA_NO_ERROR)

return res;

setup_done = 1;

}

return GA_NO_ERROR;

}

static int cuda_get_platform_count(unsigned int* platcount) {

*platcount = 1; // CUDA works on NVIDIA's GPUs

return GA_NO_ERROR;

}

static int cuda_get_device_count(unsigned int platform,

unsigned int* devcount) {

int dv;

// platform number gets ignored in CUDA implementation

GA_CHECK(setup_lib());

err = cuDeviceGetCount(&dv);

if (err != CUDA_SUCCESS)

return GA_IMPL_ERROR;

*devcount = (unsigned int)dv;

return GA_NO_ERROR;

}

cuda_context *cuda_make_ctx(CUcontext ctx, int flags) {

cuda_context *res;

void *p;

int e;

e = setup_lib();

if (e != GA_NO_ERROR)

return NULL;

res = calloc(1, sizeof(*res));

if (res == NULL)

return NULL;

res->ctx = ctx;

res->ops = &cuda_ops;

res->err = CUDA_SUCCESS;

res->refcnt = 1;

res->flags = flags;

res->enter = 0;

res->major = major;

res->minor = minor;

res->freeblocks = NULL;

if (detect_arch(ARCH_PREFIX, res->bin_id, &err)) {

goto fail_stream;

}

/* Don't add the nonblocking flags to help usage with other

libraries that may do stuff on the NULL stream */

err = cuStreamCreate(&res->s, 0);

if (err != CUDA_SUCCESS) {

goto fail_stream;

}

if (ISSET(res->flags, GA_CTX_SINGLE_STREAM)) {

res->mem_s = res->s;

} else {

/* Don't add the nonblocking flags to help usage with other

libraries that may do stuff on the NULL stream */

err = cuStreamCreate(&res->mem_s, 0);

if (err != CUDA_SUCCESS) {

goto fail_mem_stream;

}

res->kernel_cache = cache_twoq(64, 128, 64, 8, strb_eq, strb_hash,

(cache_freek_fn)strb_free,

(cache_freev_fn)cuda_freekernel);

if (res->kernel_cache == NULL)

goto fail_cache;

err = cuMemAllocHost(&p, 16);

if (err != CUDA_SUCCESS) {

goto fail_errbuf;

}

memset(p, 0, 16);

/* Need to tag for new_gpudata */

TAG_CTX(res);

res->errbuf = new_gpudata(res, (CUdeviceptr)p, 16);

if (res->errbuf == NULL) {

err = res->err;

goto fail_end;

}

res->errbuf->flags |= CUDA_MAPPED_PTR;

return res;

fail_end:

cuMemFreeHost(p);

fail_errbuf:

cache_destroy(res->kernel_cache);

fail_cache:

if (ISCLR(res->flags, GA_CTX_SINGLE_STREAM))

cuStreamDestroy(res->mem_s);

fail_mem_stream:

cuStreamDestroy(res->s);

fail_stream:

free(res);

return NULL;

}

static void deallocate(gpudata *);

static void cuda_free_ctx(cuda_context *ctx) {

gpuarray_blas_ops *blas_ops;

gpudata *next, *curr;

CUdevice dev;

ASSERT_CTX(ctx);

ctx->refcnt--;

if (ctx->refcnt == 0) {

assert(ctx->enter == 0 && "Context was active when freed!");

if (ctx->blas_handle != NULL) {

cuda_property((gpucontext *)ctx, NULL, NULL, GA_CTX_PROP_BLAS_OPS,

&blas_ops);

blas_ops->teardown((gpucontext *)ctx);

}

cuMemFreeHost((void *)ctx->errbuf->ptr);

deallocate(ctx->errbuf);

if (ISCLR(ctx->flags, GA_CTX_SINGLE_STREAM))

cuStreamDestroy(ctx->mem_s);

cuStreamDestroy(ctx->s);

/* Clear out the freelist */

for (curr = ctx->freeblocks; curr != NULL; curr = next) {

next = curr->next;

cuMemFree(curr->ptr);

deallocate(curr);

}

cache_destroy(ctx->kernel_cache);

if (!(ctx->flags & DONTFREE)) {

cuCtxPushCurrent(ctx->ctx);

cuCtxGetDevice(&dev);

cuCtxPopCurrent(NULL);

cuDevicePrimaryCtxRelease(dev);

}

CLEAR(ctx);

free(ctx);

}

CUstream cuda_get_stream(cuda_context *ctx) {

ASSERT_CTX(ctx);

return ctx->s;

}

void cuda_enter(cuda_context *ctx) {

ASSERT_CTX(ctx);

if (!ctx->enter)

cuCtxPushCurrent(ctx->ctx);

ctx->enter++;

}

void cuda_exit(cuda_context *ctx) {

ASSERT_CTX(ctx);

assert(ctx->enter > 0);

ctx->enter--;

if (!ctx->enter)

cuCtxPopCurrent(NULL);

}

static gpudata *new_gpudata(cuda_context *ctx, CUdeviceptr ptr, size_t size) {

gpudata *res;

int fl = CU_EVENT_DISABLE_TIMING;

res = malloc(sizeof(*res));

if (res == NULL) return NULL;

res->refcnt = 0;

res->sz = size;

res->flags = 0;

res->ls = NULL;

cuda_enter(ctx);

if (ctx->flags & GA_CTX_MULTI_THREAD)

fl |= CU_EVENT_BLOCKING_SYNC;

ctx->err = cuEventCreate(&res->rev, fl);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

free(res);

return NULL;

}

ctx->err = cuEventCreate(&res->wev, fl);

if (ctx->err != CUDA_SUCCESS) {

cuEventDestroy(res->rev);

cuda_exit(ctx);

free(res);

return NULL;

}

cuda_exit(ctx);

res->ptr = ptr;

res->next = NULL;

res->ctx = ctx;

TAG_BUF(res);

return res;

}

gpudata *cuda_make_buf(cuda_context *ctx, CUdeviceptr p, size_t sz) {

gpudata *res = new_gpudata(ctx, p, sz);

if (res == NULL) return NULL;

res->refcnt = 1;

res->flags |= DONTFREE;

res->ctx->refcnt++;

return res;

}

size_t cuda_get_sz(gpudata *g) { ASSERT_BUF(g); return g->sz; }

#define FAIL(v, e) { if (ret) *ret = e; return v; }

#define CHKFAIL(v) if (err != CUDA_SUCCESS) FAIL(v, GA_IMPL_ERROR)

static const char CUDA_PREAMBLE[] =

"#define local_barrier() __syncthreads()\n"

"#define WITHIN_KERNEL extern \"C\" __device__\n"

"#define KERNEL extern \"C\" __global__\n"

"#define GLOBAL_MEM /* empty */\n"

"#define LOCAL_MEM __shared__\n"

"#define LOCAL_MEM_ARG /* empty */\n"

"#define REQD_WG_SIZE(X,Y,Z) __launch_bounds__(X*Y, Z)\n"

"#ifdef NAN\n"

"#undef NAN\n"

"#endif\n"

"#define NAN __int_as_float(0x7fffffff)\n"

"#define LID_0 threadIdx.x\n"

"#define LID_1 threadIdx.y\n"

"#define LID_2 threadIdx.z\n"

"#define LDIM_0 blockDim.x\n"

"#define LDIM_1 blockDim.y\n"

"#define LDIM_2 blockDim.z\n"

"#define GID_0 blockIdx.x\n"

"#define GID_1 blockIdx.y\n"

"#define GID_2 blockIdx.z\n"

"#define GDIM_0 gridDim.x\n"

"#define GDIM_1 gridDim.y\n"

"#define GDIM_2 gridDim.z\n"

"#define ga_bool unsigned char\n"

"#define ga_byte signed char\n"

"#define ga_ubyte unsigned char\n"

"#define ga_short short\n"

"#define ga_ushort unsigned short\n"

"#define ga_int int\n"

"#define ga_uint unsigned int\n"

"#define ga_long long long\n"

"#define ga_ulong unsigned long long\n"

"#define ga_float float\n"

"#define ga_double double\n"

"#define ga_half ga_ushort\n"

"#define ga_size size_t\n"

"#define ga_ssize ptrdiff_t\n"

"#define load_half(p) __half2float(*(p))\n"

"#define store_half(p, v) (*(p) = __float2half_rn(v))\n"

"#define GA_DECL_SHARED_PARAM(type, name)\n"

"#define GA_DECL_SHARED_BODY(type, name) extern __shared__ type name[];\n"

"#define GA_WARP_SIZE warpSize\n"

"#line 1\n";

/* XXX: add complex, quads, longlong */

/* XXX: add vector types */

static cuda_context *do_init(CUdevice dev, int flags, int *ret) {

cuda_context *res;

CUcontext ctx;

unsigned int fl = CU_CTX_SCHED_AUTO;

unsigned int cur_fl;

int act;

int i;

CHKFAIL(NULL);

if (flags & GA_CTX_SINGLE_THREAD)

fl = CU_CTX_SCHED_SPIN;

if (flags & GA_CTX_MULTI_THREAD)

fl = CU_CTX_SCHED_YIELD;

err = cuDeviceGetAttribute(&i, CU_DEVICE_ATTRIBUTE_UNIFIED_ADDRESSING, dev);

CHKFAIL(NULL);

if (i != 1)

FAIL(NULL, GA_UNSUPPORTED_ERROR);

err = cuDevicePrimaryCtxGetState(dev, &cur_fl, &act);

CHKFAIL(NULL);

if (act == 1) {

if ((cur_fl & fl) != fl)

FAIL(NULL, GA_INVALID_ERROR);

} else {

err = cuDevicePrimaryCtxSetFlags(dev, fl);

CHKFAIL(NULL);

}

err = cuDevicePrimaryCtxRetain(&ctx, dev);

CHKFAIL(NULL);

err = cuCtxPushCurrent(ctx);

CHKFAIL(NULL);

res = cuda_make_ctx(ctx, flags);

if (res == NULL) {

cuDevicePrimaryCtxRelease(dev);

FAIL(NULL, GA_IMPL_ERROR);

}

/* Don't leave the context on the thread stack */

cuCtxPopCurrent(NULL);

return res;

}

static gpucontext *cuda_init(int ord, int flags, int *ret) {

CUdevice dev;

cuda_context *res;

int r;

r = setup_lib();

if (r != GA_NO_ERROR) {

FAIL(NULL, r);

}

if (ord == -1) {

int i, c;

err = cuDeviceGetCount(&c);

CHKFAIL(NULL);

for (i = 0; i < c; i++) {

err = cuDeviceGet(&dev, i);

CHKFAIL(NULL);

res = do_init(dev, flags, NULL);

if (res != NULL)

return (gpucontext *)res;

}

FAIL(NULL, GA_NODEV_ERROR);

} else {

err = cuDeviceGet(&dev, ord);

CHKFAIL(NULL);

return (gpucontext *)do_init(dev, flags, ret);

}

static void cuda_deinit(gpucontext *c) {

cuda_free_ctx((cuda_context *)c);

}

* Find the block in the free list that is the best fit for the size

* we want, which means the smallest that can still fit the size.

static void find_best(cuda_context *ctx, gpudata **best, gpudata **prev,

size_t size) {

gpudata *temp, *tempPrev = NULL;

*best = NULL;

for (temp = ctx->freeblocks; temp; temp = temp->next) {

if (temp->sz >= size && (!*best || temp->sz < (*best)->sz)) {

*best = temp;

*prev = tempPrev;

}

tempPrev = temp;

}

static size_t largest_size(cuda_context *ctx) {

gpudata *temp;

size_t sz, dummy;

cuda_enter(ctx);

ctx->err = cuMemGetInfo(&sz, &dummy);

cuda_exit(ctx);

/* We guess that we can allocate at least a quarter of the free size

in a single block. This might be wrong though. */

sz /= 4;

for (temp = ctx->freeblocks; temp; temp = temp->next) {

if (temp->sz > sz) sz = temp->sz;

}

return sz;

}

* Allocate a new block and place in on the freelist. Will allocate

* the bigger of the requested size and BLOCK_SIZE to avoid allocating

* multiple small blocks.

static int allocate(cuda_context *ctx, gpudata **res, gpudata **prev,

size_t size) {

CUdeviceptr ptr;

gpudata *next;

*prev = NULL;

if (!(ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE))

if (size < BLOCK_SIZE) size = BLOCK_SIZE;

cuda_enter(ctx);

ctx->err = cuMemAlloc(&ptr, size);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

*res = new_gpudata(ctx, ptr, size);

cuda_exit(ctx);

if (*res == NULL) {

cuMemFree(ptr);

return GA_MEMORY_ERROR;

}

(*res)->flags |= CUDA_HEAD_ALLOC;

/* Now that the block is allocated, enter it in the freelist */

next = ctx->freeblocks;

for (; next && next->ptr < (*res)->ptr; next = next->next) {

*prev = next;

}

(*res)->next = next;

if (*prev)

(*prev)->next = *res;

else

ctx->freeblocks = *res;

return GA_NO_ERROR;

}

* Extract the `curr` block from the freelist, possibly splitting it

* if it's too big for the requested size. The remaining block will

* stay on the freelist if there is a split. `prev` is only to

* facilitate the extraction so we don't have to go through the list

* again.

static int extract(gpudata *curr, gpudata *prev, size_t size) {

gpudata *next, *split;

size_t remaining = curr->sz - size;

if (remaining < FRAG_SIZE) {

/* No need to split, the remaining block would be too small */

next = curr->next;

} else {

split = new_gpudata(curr->ctx, curr->ptr + size, remaining);

if (split == NULL)

return GA_MEMORY_ERROR;

/* Make sure the chain keeps going */

split->next = curr->next;

curr->next = NULL;

/* Make sure we don't start using the split buffer too soon */

cuda_records(split, CUDA_WAIT_ALL, curr->ls);

next = split;

curr->sz = size;

}

if (prev != NULL)

prev->next = next;

else

curr->ctx->freeblocks = next;

return GA_NO_ERROR;

}

static void cuda_free(gpudata *);

static int cuda_write(gpudata *dst, size_t dstoff, const void *src,

size_t sz);

static inline size_t roundup(size_t s, size_t m) {

return ((s + (m - 1)) / m) * m;

}

static gpudata *cuda_alloc(gpucontext *c, size_t size, void *data, int flags,

int *ret) {

gpudata *res = NULL, *prev = NULL;

cuda_context *ctx = (cuda_context *)c;

size_t asize;

int err;

if ((flags & GA_BUFFER_INIT) && data == NULL) FAIL(NULL, GA_VALUE_ERROR);

if ((flags & (GA_BUFFER_READ_ONLY|GA_BUFFER_WRITE_ONLY)) ==

(GA_BUFFER_READ_ONLY|GA_BUFFER_WRITE_ONLY)) FAIL(NULL, GA_VALUE_ERROR);

/* TODO: figure out how to make this work */

if (flags & GA_BUFFER_HOST) FAIL(NULL, GA_DEVSUP_ERROR);

/* We don't want to manage really small allocations so we round up

* to a multiple of FRAG_SIZE. This also ensures that if we split a

* block, the next block starts properly aligned for any data type.

if (!(ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE)) {

asize = roundup(size, FRAG_SIZE);

find_best(ctx, &res, &prev, asize);

} else {

asize = size;

}

if (res == NULL) {

err = allocate(ctx, &res, &prev, asize);

if (err != GA_NO_ERROR)

FAIL(NULL, err);

}

err = extract(res, prev, asize);

if (err != GA_NO_ERROR)

FAIL(NULL, err);

/* It's out of the freelist, so add a ref */

res->ctx->refcnt++;

/* We consider this buffer allocated and ready to go */

res->refcnt = 1;

if (flags & GA_BUFFER_INIT) {

err = cuda_write(res, 0, data, size);

if (err != GA_NO_ERROR) {

cuda_free(res);

FAIL(NULL, err);

}

return res;

}

int cuda_get_ipc_handle(gpudata *d, GpuArrayIpcMemHandle *h) {

ASSERT_BUF(d);

cuda_enter(d->ctx);

CUDA_EXIT_ON_ERROR(d->ctx,

cuIpcGetMemHandle((CUipcMemHandle *)h, d->ptr));

cuda_exit(d->ctx);

return GA_NO_ERROR;

}

gpudata *cuda_open_ipc_handle(gpucontext *c, GpuArrayIpcMemHandle *h, size_t sz) {

CUdeviceptr p;

cuda_context *ctx = (cuda_context *)c;

gpudata *d = NULL;

cuda_enter(ctx);

ctx->err = cuIpcOpenMemHandle(&p, *((CUipcMemHandle *)h),

CU_IPC_MEM_LAZY_ENABLE_PEER_ACCESS);

if (ctx->err == CUDA_SUCCESS) {

d = cuda_make_buf(ctx, p, sz);

if (d != NULL)

d->flags |= CUDA_IPC_MEMORY;

}

cuda_exit(ctx);

return d;

}

static void cuda_retain(gpudata *d) {

ASSERT_BUF(d);

d->refcnt++;

}

static void deallocate(gpudata *d) {

cuda_enter(d->ctx);

cuEventDestroy(d->rev);

cuEventDestroy(d->wev);

cuda_exit(d->ctx);

CLEAR(d);

free(d);

}

static void cuda_free(gpudata *d) {

/* We ignore errors on free */

ASSERT_BUF(d);

d->refcnt--;

if (d->refcnt == 0) {

/* Keep a reference to the context since we deallocate the gpudata

* object */

cuda_context *ctx = d->ctx;

if (d->flags & DONTFREE) {

/* This is the path for "external" buffers */

deallocate(d);

} else if (d->flags & CUDA_IPC_MEMORY) {

cuIpcCloseMemHandle(d->ptr);

deallocate(d);

} else if (ctx->flags & GA_CTX_DISABLE_ALLOCATION_CACHE) {

/* Just free the pointer */

cuMemFree(d->ptr);

deallocate(d);

} else {

/* Find the position in the freelist. Freelist is kept in order

of allocation address */

gpudata *next = d->ctx->freeblocks, *prev = NULL;

for (; next && next->ptr < d->ptr; next = next->next) {

prev = next;

}

next = prev != NULL ? prev->next : d->ctx->freeblocks;

/* See if we can merge the block with the previous one */

if (!(d->flags & CUDA_HEAD_ALLOC) &&

prev != NULL && prev->ptr + prev->sz == d->ptr) {

prev->sz = prev->sz + d->sz;

cuda_waits(d, CUDA_WAIT_ALL, prev->ls);

cuda_records(prev, CUDA_WAIT_ALL, prev->ls);

deallocate(d);

d = prev;

} else if (prev != NULL) {

prev->next = d;

} else {

d->ctx->freeblocks = d;

}

/* See if we can merge with next */

if (next && !(next->flags & CUDA_HEAD_ALLOC) &&

d->ptr + d->sz == next->ptr) {

d->sz = d->sz + next->sz;

d->next = next->next;

cuda_wait(next, CUDA_WAIT_ALL);

cuda_record(d, CUDA_WAIT_ALL);

deallocate(next);

} else {

d->next = next;

}

/* We keep this at the end since the freed buffer could be the

* last reference to the context and therefore clearing the

* reference could trigger the freeing if the whole context

* including the freelist, which we manipulate. */

cuda_free_ctx(ctx);

}

static int cuda_share(gpudata *a, gpudata *b, int *ret) {

ASSERT_BUF(a);

ASSERT_BUF(b);

return (a->ctx == b->ctx && a->sz != 0 && b->sz != 0 &&

((a->ptr <= b->ptr && a->ptr + a->sz > b->ptr) ||

(b->ptr <= a->ptr && b->ptr + b->sz > a->ptr)));

}

static int cuda_waits(gpudata *a, int flags, CUstream s) {

ASSERT_BUF(a);

/* Never skip the wait if CUDA_WAIT_FORCE */

if (ISCLR(flags, CUDA_WAIT_FORCE)) {

if (ISSET(a->ctx->flags, GA_CTX_SINGLE_STREAM))

return GA_NO_ERROR;

/* If the last stream to touch this buffer is the same, we don't

* need to wait for anything. */

if (a->ls == s)

return GA_NO_ERROR;

}

cuda_enter(a->ctx);

/* We wait for writes that happened before since multiple reads at

* the same time are fine */

if (ISSET(flags, CUDA_WAIT_READ) || ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuStreamWaitEvent(s, a->wev, 0));

/* Make sure to not disturb previous reads */

if (ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuStreamWaitEvent(s, a->rev, 0));

cuda_exit(a->ctx);

return GA_NO_ERROR;

}

int cuda_wait(gpudata *a, int flags) {

return cuda_waits(a, flags, a->ctx->s);

}

static int cuda_records(gpudata *a, int flags, CUstream s) {

ASSERT_BUF(a);

if (ISCLR(flags, CUDA_WAIT_FORCE) &&

ISSET(a->ctx->flags, GA_CTX_SINGLE_STREAM))

return GA_NO_ERROR;

cuda_enter(a->ctx);

if (ISSET(flags, CUDA_WAIT_READ))

CUDA_EXIT_ON_ERROR(a->ctx, cuEventRecord(a->rev, s));

if (ISSET(flags, CUDA_WAIT_WRITE))

CUDA_EXIT_ON_ERROR(a->ctx, cuEventRecord(a->wev, s));

cuda_exit(a->ctx);

a->ls = s;

return GA_NO_ERROR;

}

int cuda_record(gpudata *a, int flags) {

return cuda_records(a, flags, a->ctx->s);

}

static int cuda_move(gpudata *dst, size_t dstoff, gpudata *src,

size_t srcoff, size_t sz) {

cuda_context *ctx = dst->ctx;

int res = GA_NO_ERROR;

ASSERT_BUF(dst);

ASSERT_BUF(src);

if (src->ctx != dst->ctx) return GA_VALUE_ERROR;

if (sz == 0) return GA_NO_ERROR;

if ((dst->sz - dstoff) < sz || (src->sz - srcoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(src, CUDA_WAIT_READ));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(dst, CUDA_WAIT_WRITE));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyDtoDAsync(dst->ptr + dstoff, src->ptr + srcoff, sz, ctx->s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(src, CUDA_WAIT_READ));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(dst, CUDA_WAIT_WRITE));

cuda_exit(ctx);

return res;

}

static int cuda_read(void *dst, gpudata *src, size_t srcoff, size_t sz) {

cuda_context *ctx = src->ctx;

ASSERT_BUF(src);

if (sz == 0) return GA_NO_ERROR;

if ((src->sz - srcoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

if (src->flags & CUDA_MAPPED_PTR) {

if (ISSET(ctx->flags, GA_CTX_SINGLE_STREAM))

ctx->err = cuStreamSynchronize(ctx->s);

else

ctx->err = cuEventSynchronize(src->wev);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

memcpy(dst, (void *)(src->ptr + srcoff), sz);

} else {

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_waits(src, CUDA_WAIT_READ, ctx->mem_s));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyDtoHAsync(dst, src->ptr + srcoff, sz, ctx->mem_s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_records(src, CUDA_WAIT_READ, ctx->mem_s));

}

cuda_exit(ctx);

return GA_NO_ERROR;

}

static int cuda_write(gpudata *dst, size_t dstoff, const void *src,

size_t sz) {

cuda_context *ctx = dst->ctx;

ASSERT_BUF(dst);

if (sz == 0) return GA_NO_ERROR;

if ((dst->sz - dstoff) < sz)

return GA_VALUE_ERROR;

cuda_enter(ctx);

if (dst->flags & CUDA_MAPPED_PTR) {

if (ISSET(ctx->flags, GA_CTX_SINGLE_STREAM))

ctx->err = cuStreamSynchronize(ctx->s);

else

ctx->err = cuEventSynchronize(dst->rev);

if (ctx->err != CUDA_SUCCESS) {

cuda_exit(ctx);

return GA_IMPL_ERROR;

}

memcpy((void *)(dst->ptr + dstoff), src, sz);

} else {

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_waits(dst, CUDA_WAIT_WRITE, ctx->mem_s));

CUDA_EXIT_ON_ERROR(ctx,

cuMemcpyHtoDAsync(dst->ptr + dstoff, src, sz, ctx->mem_s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_records(dst, CUDA_WAIT_WRITE, ctx->mem_s));

}

cuda_exit(ctx);

return GA_NO_ERROR;

}

static int cuda_memset(gpudata *dst, size_t dstoff, int data) {

cuda_context *ctx = dst->ctx;

ASSERT_BUF(dst);

if ((dst->sz - dstoff) == 0) return GA_NO_ERROR;

cuda_enter(ctx);

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_wait(dst, CUDA_WAIT_WRITE));

CUDA_EXIT_ON_ERROR(ctx,

cuMemsetD8Async(dst->ptr + dstoff, data, dst->sz - dstoff, ctx->s));

GA_CUDA_EXIT_ON_ERROR(ctx,

cuda_record(dst, CUDA_WAIT_WRITE));

cuda_exit(ctx);

return GA_NO_ERROR;

}

static CUresult get_cc(CUdevice dev, int *maj, int *min) {

CUresult lerr;

lerr = cuDeviceGetAttribute(maj,

CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR,

dev);

if (lerr != CUDA_SUCCESS)

return lerr;

return cuDeviceGetAttribute(min,

CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR,

dev);

}

static int detect_arch(const char *prefix, char *ret, CUresult *err) {

CUdevice dev;

int major, minor;

int res;

size_t sz = strlen(prefix) + 3;

*err = cuCtxGetDevice(&dev);

if (*err != CUDA_SUCCESS) return GA_IMPL_ERROR;

*err = get_cc(dev, &major, &minor);

if (*err != CUDA_SUCCESS) return GA_IMPL_ERROR;

res = snprintf(ret, sz, "%s%d%d", prefix, major, minor);

if (res == -1 || res > (ssize_t)sz) return GA_UNSUPPORTED_ERROR;

return GA_NO_ERROR;

}

static void *call_compiler(const char *src, size_t len, const char *arch_arg,

size_t *bin_len, char **log, size_t *log_len,

int *ret) {

nvrtcProgram prog;

void *buf = NULL;

size_t buflen;

const char *opts[4] = {

"-arch", ""

, "-G", "-lineinfo"

};

nvrtcResult err, err2;

opts[1] = arch_arg;

err = nvrtcCreateProgram(&prog, src, NULL, 0, NULL, NULL);

if (err != NVRTC_SUCCESS) FAIL(NULL, GA_SYS_ERROR);

err = nvrtcCompileProgram(prog,

#ifdef DEBUG

#else

#endif

opts);

if (log != NULL) {

err2 = nvrtcGetProgramLogSize(prog, &buflen);

if (err2 != NVRTC_SUCCESS) goto end2;

buf = malloc(buflen);

if (buf == NULL) goto end2;

err2 = nvrtcGetProgramLog(prog, (char *)buf);

if (err2 != NVRTC_SUCCESS) goto end2;

if (log_len != NULL) *log_len = buflen;

*log = (char *)buf;

buf = NULL;

}

end2:

if (err != NVRTC_SUCCESS) goto end;

err = nvrtcGetPTXSize(prog, &buflen);

if (err != NVRTC_SUCCESS) goto end;

buf = malloc(buflen);

if (buf == NULL) {

nvrtcDestroyProgram(&prog);

FAIL(NULL, GA_MEMORY_ERROR);

}

err = nvrtcGetPTX(prog, (char *)buf);

if (err != NVRTC_SUCCESS) goto end;

*bin_len = buflen;

end:

nvrtcDestroyProgram(&prog);

if (err != NVRTC_SUCCESS) {

free(buf);

FAIL(NULL, GA_SYS_ERROR);

}

return buf;

}

static void _cuda_freekernel(gpukernel *k) {

k->refcnt--;

if (k->refcnt == 0) {

if (k->ctx != NULL) {

cuda_enter(k->ctx);

cuModuleUnload(k->m);

cuda_exit(k->ctx);

cuda_free_ctx(k->ctx);

}

CLEAR(k);

free(k->args);

free(k->bin);

free(k->types);

free(k);

}

static gpukernel *cuda_newkernel(gpucontext *c, unsigned int count,

const char **strings, const size_t *lengths,

const char *fname, unsigned int argcount,

const int *types, int flags, int *ret,

char **err_str) {

cuda_context *ctx = (cuda_context *)c;

strb sb = STRB_STATIC_INIT;

strb *psb;

char *bin, *log = NULL;

gpukernel *res;

size_t bin_len = 0, log_len = 0;

View remainder of file in raw view

libgpuarray/src/gpuarray_buffer_cuda.c at master · jrao1/libgpuarray · GitHub

Expand file tree

Breadcrumbs

gpuarray_buffer_cuda.c

Latest commit

History

Breadcrumbs

gpuarray_buffer_cuda.c

File metadata and controls