Av: Didrik Pemmer Aalen
Et Harshadtall er definert som et positivt tall som er delelig med summen av alle sifferne i tallet. I noen tilfeller er summen av alle sifferne i tallet et primtall. Disse kaller vi Harshadprimtall.
1729 er et Harshadtall fordi 1 + 7 + 2 + 9 = 19 og 1729 % 19 = 0. Dette er også et Harshadprimtall, fordi 19 er et primtall.
1730 er ikke et Harshadtall fordi 1 + 7 + 3 + 0 = 11 og 1730 % 11 = 3.
Hvor mange tall fra 1 til og med 98765432 er Harshadprimtall?
Jeg har vært inne på tanken selv. Rekka i seg selv er jo egentlig ganske rett frem, men 10 nøstede for-løkker er jo litt voldsomt. Kunne kanskje rullet ut 2 og 2 siffer, og nøste 5 slike rekker?
Min strategi er jo egentlig å droppe løkker og loop-variabler mest mulig slik at man kan slippe branchingen det fører til. Unntaket er addToDigit som jeg ikke fikk mer effektiv ved å unrolle, men denne kalles uansett bare 1 gang for hvert 100 tall.
@michaelo
Hmm, fikk akkurat en ny idé... Hvis jeg kan splitte arbeidet i litt penere aligned chunker før trådene setter i gang blir jo egentlig alt mye lettere. Da kan jeg kanskje bake inn noen antakelser om modulo 6 også... Fristende! Men det må bli en annen dag.
@michaelo @bobcat4 @DarioSucic
Jeg har tenkt litt... i rettferdighetens navn må man kunne argumentere for at når man sammenlikner benchmarks kjørt på et skjermkort til 9000 kroner så må det være rimelig å sammenlikne dem med tilsvarende benchmarks kjørt på en prosessor til samme pris.
Så, jeg fant en kollega med en 3950X og tigget meg til litt kjøretid på hans private boks. Her er resultatet:
Lite, ST: 34,5 ms
Lite, MT (precalc): 3,78ms
Lite, MT: 3,74ms
Stort, ST: 712ms
Stort. MT (precalc): 57ms
Stort. MT: 57ms
@terjew cc: @DarioSucic @bobcat4
Hehe. Bra bra. Gyldig argument det!
Så er da de neste spørsmålene:
?
Fikk ned CUDA fra 5.4ms til 4.1ms ved å bytte fra Atomic Add til Parallell blokk-reduksjon slik det er beskrevet her.
Tror det er mye mer å hente, men det er i alle fall en start. 😅
#include <stdlib.h>
#include <stdint.h>
#include <stdio.h>
#include <math.h>
#define u32 uint32_t
#define u8 uint8_t
#define f32 float
#define f64 double
#define BLOCK_SIZE 128
__device__ void warp_reduce(volatile u32 *block_data, u32 tid)
{
block_data[tid] += block_data[tid + 32];
block_data[tid] += block_data[tid + 16];
block_data[tid] += block_data[tid + 8];
block_data[tid] += block_data[tid + 4];
block_data[tid] += block_data[tid + 2];
block_data[tid] += block_data[tid + 1];
}
__device__ u32 compute_single(u32 k, u8 *prime)
{
const u32 pow_table[9] = {1, 10, 100, 1000, 10000, 100000, 1000000, 10000000, 100000000};
u32 digit_sum = 0;
for (u32 i = 0; i < 9; i++)
digit_sum += (k / pow_table[i]) % 10;
return (k % digit_sum == 0) & prime[digit_sum];
}
__global__ void kernel(u32 *block_counts, u8 *prime, u32 n)
{
extern __shared__ u32 block_data[];
u32 tid = threadIdx.x;
u32 idx = blockIdx.x * blockDim.x + tid;
if (idx > n) block_data[tid] = 0;
else block_data[tid] = compute_single(idx + 1, prime);
__syncthreads();
for(u32 s = blockDim.x / 2; s > 32; s >>= 1)
{
if(tid < s) block_data[tid] += block_data[tid + s];
__syncthreads();
}
if(tid < 32) warp_reduce(block_data, tid);
if(tid == 0) block_counts[blockIdx.x] = block_data[0];
}
u8 *sieve(u32 n)
{
u32 lim = n + 1;
u8 *prime = (u8 *)malloc(lim);
memset(prime, 1, lim);
prime[0] = 0;
prime[1] = 0;
for (u32 k = 2; k < lim; k++)
{
if (!prime[k])
continue;
for (u32 i = k * k; i < lim; i += k)
prime[i] = 0;
}
return prime;
}
int main()
{
const u32 n = 98765432;
const u32 prime_limit = 9 * ceil(log10(n));
cudaEvent_t st, et;
cudaEventCreate(&st);
cudaEventCreate(&et);
u8 *prime = sieve(prime_limit);
u8 *prime_device;
cudaMalloc(&prime_device, prime_limit);
cudaMemcpy(prime_device, prime, prime_limit, cudaMemcpyHostToDevice);
u32 num_blocks = (n + BLOCK_SIZE - 1) / BLOCK_SIZE;
u32 shared_mem_size = BLOCK_SIZE * sizeof(u32);
u32 *counts = (u32 *) malloc(num_blocks * sizeof(u32));
u32 *counts_device;
cudaMalloc(&counts_device, num_blocks * sizeof(u32));
cudaEventRecord(st);
kernel<<<num_blocks, BLOCK_SIZE, shared_mem_size>>>(counts_device, prime_device, n);
cudaEventRecord(et);
cudaEventSynchronize(et);
float ms = 0;
cudaEventElapsedTime(&ms, st, et);
printf("Kernel execution: %.3f ms\n", ms);
cudaMemcpy((void *)counts, counts_device, num_blocks * sizeof(u32), cudaMemcpyDeviceToHost);
u32 total_count = 0;
for(u32 i = 0; i < num_blocks; i++)
total_count += counts[i];
printf("Count: %d\n", total_count);
cudaFree(counts_device);
cudaFree(prime_device);
free(prime);
return 0;
}
Fikk ned CUDA fra 5.4ms til 4.1ms ved å bytte fra Atomic Add til Parallell blokk-reduksjon slik det er beskrevet her.
Tror det er mye mer å hente, men det er i alle fall en start. 😅
Nais! 👏
Tar du en benchmark på stort sett også?
- Varer også kodekalenderjula helt til påske?
Minst! Og da må den vel avløses av noen kodepåskenøtter?
Tar du en benchmark på stort sett også?
Endte opp med 85ms på det store datasettet, men egentlig jukser jeg litt siden mesteparten av summeringen nå er flyttet over til CPU utenfor kernelen. Om jeg får tid kan jeg se om jeg ikke får slengt inn litt bedre timing slik at resultatet er sammenlignbart med resten av tidene deres.
Tar du en benchmark på stort sett også?
Endte opp med
85mspå det store datasettet, men egentlig jukser jeg litt siden mesteparten av summeringen nå er flyttet over til CPU utenfor kernelen. Om jeg får tid kan jeg se om jeg ikke får slengt inn litt bedre timing slik at resultatet er sammenlignbart med resten av tidene deres.
Tenkte på det samme når jeg så over koden. Du gjemmer bort et par ting i benchmarken, deriblant den siste summeringen men også memcpy til og fra CUDA-devicen. Tenker kanskje det hadde vært riktigere å la programmet loope f.eks 1000 ganger på lite datasett, og måle total tid brukt på dette inkludert de punktene, men uten å ta med opprettelse og teardown av CUDA-konteksten. Det utgjør vel ikke all verden, men når vi er nede på noen få millisekunder for hele kjøringen begynner alle sånne småting å utgjøre en relevant del av tiden.
Evt time hele programmet med 1000 interne kjøringer fra shellet slik vi gjør for de andre variantene. Da får du jo med de 200 ms på opprettelsen av CUDA-kontekst etc, men dette fordeles jo utover de 1000 kjøringene så det skal vel ikke utgjøre mer enn 0,2ms per kjøring.
@terjew
Hehe. Kjenner følelsen!
Den var fin! 👍
Jeg nevnte forøvrig i en tidligere post at jeg hadde noen spor på å utlede rekka av summer direkte, men uansett hvor lur og/eller enkel jeg har forsøkt å være så ender det opp i for mange løkker til å bli ordentlig effektivt (det ble fort noe ala kjøretidskompleksiteten til posten til @DarioSucic lenger opp her) selv om jeg jobbet direkte på summen og rullet ut det dypeste nivået. Fikk den vel ned mot 300ms st, så ikke noe å hoppe i taket av.